1 /* { dg-do run { target openacc_nvidia_accel_selected } }
2 { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
3 /* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
11 #pragma acc loop worker
12 for (v1
= 0; v1
< 20; v1
+= 2)
18 /* Todo: Boths bar.syncs can be removed.
19 Atm we generate this dead code inbetween forked and joining:
21 mov.u32 %r28, %ntid.y;
23 add.u32 %r30, %r29, %r29;
24 setp.gt.s32 %r31, %r30, 19;
26 add.u32 %r25, %r28, %r28;
29 add.u32 %r24, %r24, %r25;
30 setp.le.s32 %r33, %r24, 19;
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" } } */