1 /* { dg-do run { target openacc_nvidia_accel_selected } } */
2 /* This code uses nvptx inline assembly guarded with acc_on_device, which is
3 not optimized away at -O0, and then confuses the target assembler.
4 { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
12 (acc_on_device (acc_device_nvidia) \
14 __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (__r)); \
18 test_static(int *a
, int num_gangs
, int sarg
)
25 for (i
= 0; i
< N
/ sarg
; i
++)
26 for (j
= 0; j
< sarg
; j
++)
27 assert (a
[i
*sarg
+j
] == i
% num_gangs
);
31 test_nonstatic(int *a
, int gangs
)
35 for (i
= 0; i
< N
; i
+=gangs
)
36 for (j
= 0; j
< gangs
; j
++)
37 assert (a
[i
+j
] == i
/gangs
);
46 #pragma acc parallel loop gang (static:*) num_gangs (10)
47 for (i
= 0; i
< 100; i
++)
50 test_nonstatic (a
, 10);
52 #pragma acc parallel loop gang (static:1) num_gangs (10)
53 for (i
= 0; i
< 100; i
++)
56 test_static (a
, 10, 1);
58 #pragma acc parallel loop gang (static:2) num_gangs (10)
59 for (i
= 0; i
< 100; i
++)
62 test_static (a
, 10, 2);
64 #pragma acc parallel loop gang (static:5) num_gangs (10)
65 for (i
= 0; i
< 100; i
++)
68 test_static (a
, 10, 5);
70 #pragma acc parallel loop gang (static:20) num_gangs (10)
71 for (i
= 0; i
< 100; i
++)
74 test_static (a
, 10, 20);
76 /* Non-static gang. */
77 #pragma acc parallel loop gang num_gangs (10)
78 for (i
= 0; i
< 100; i
++)
81 test_nonstatic (a
, 10);
83 /* Static arguments with a variable expression. */
86 #pragma acc parallel loop gang (static:0+x) num_gangs (10)
87 for (i
= 0; i
< 100; i
++)
90 test_static (a
, 10, 20);
93 #pragma acc parallel loop gang (static:x) num_gangs (10)
94 for (i
= 0; i
< 100; i
++)
97 test_static (a
, 10, 20);