1 /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
2 aspects of that functionality. */
3
4 #include <stdio.h>
5 #include <openacc.h>
6 #include <gomp-constants.h>
7
8 #define N (32*32*32+17)
9
10 #pragma acc routine gang
11 void __attribute__ ((noinline)) gang (int ary[N])
12 /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-1 } */
13 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-2 } */
14 {
15 #pragma acc loop gang
16 for (unsigned ix = 0; ix < N; ix++)
17 {
18 if (acc_on_device (acc_device_not_host))
19 {
20 int g, w, v;
21
22 g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
23 w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
24 v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
25 ary[ix] = (g << 16) | (w << 8) | v;
26 }
27 else
28 ary[ix] = ix;
29 }
30 }
31
32 int main ()
33 {
34 int ary[N];
35 int ix;
36 int exit = 0;
37 int ondev = 0;
38
39 for (ix = 0; ix < N;ix++)
40 ary[ix] = -1;
41
42 #pragma acc parallel num_gangs(32) copy(ary) copy(ondev)
43 {
44 ondev = acc_on_device (acc_device_not_host);
45 gang (ary);
46 }
47
48 for (ix = 0; ix < N; ix++)
49 {
50 int expected = ix;
51 if(ondev)
52 {
53 int g = ix / ((N + 31) / 32);
54 int w = 0;
55 int v = 0;
56
57 expected = (g << 16) | (w << 8) | v;
58 }
59
60 if (ary[ix] != expected)
61 {
62 exit = 1;
63 printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
64 }
65 }
66
67 return exit;
68 }