| /* { dg-do run { target openacc_nvidia_accel_selected } } |
| { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ |
| /* { dg-additional-options "-foffload=-fdump-rtl-mach" } */ |
| |
| int |
| main (void) |
| { |
| int v1; |
| |
| #pragma acc parallel |
| #pragma acc loop worker |
| for (v1 = 0; v1 < 20; v1 += 2) |
| ; |
| |
| return 0; |
| } |
| |
| /* Todo: Boths bar.syncs can be removed. |
| Atm we generate this dead code inbetween forked and joining: |
| |
| mov.u32 %r28, %ntid.y; |
| mov.u32 %r29, %tid.y; |
| add.u32 %r30, %r29, %r29; |
| setp.gt.s32 %r31, %r30, 19; |
| @%r31 bra $L2; |
| add.u32 %r25, %r28, %r28; |
| mov.u32 %r24, %r30; |
| $L3: |
| add.u32 %r24, %r24, %r25; |
| setp.le.s32 %r33, %r24, 19; |
| @%r33 bra $L3; |
| $L2: |
| |
| so the loop is not recognized as empty loop (which we detect by seeing if |
| joining immediately follows forked). */ |
| /* { dg-final { scan-offload-rtl-dump-times "nvptx_barsync" 2 "mach" } } */ |
