blob: 84b9c01443e50df451f6dbf50f6a42ae90bfb1cb (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
|
/* { 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" } } */
|
