1  /* { dg-do run { target openacc_nvidia_accel_selected } }
       2     { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
       3  /* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
       4  
       5  int
       6  main (void)
       7  {
       8    int v1;
       9  
      10    #pragma acc parallel
      11    #pragma acc loop worker
      12    for (v1 = 0; v1 < 20; v1 += 2)
      13      ;
      14  
      15    return 0;
      16  }
      17  
      18  /* Todo: Boths bar.syncs can be removed.
      19     Atm we generate this dead code inbetween forked and joining:
      20  
      21                       mov.u32 %r28, %ntid.y;
      22                       mov.u32 %r29, %tid.y;
      23                       add.u32 %r30, %r29, %r29;
      24                       setp.gt.s32     %r31, %r30, 19;
      25               @%r31   bra     $L2;
      26                       add.u32 %r25, %r28, %r28;
      27                       mov.u32 %r24, %r30;
      28       $L3:
      29                       add.u32 %r24, %r24, %r25;
      30                       setp.le.s32     %r33, %r24, 19;
      31               @%r33   bra     $L3;
      32       $L2:
      33  
      34     so the loop is not recognized as empty loop (which we detect by seeing if
      35     joining immediately follows forked).  */
      36  /* { dg-final { scan-offload-rtl-dump-times "nvptx_barsync" 2 "mach" } } */