libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c

   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" } } */