1 #include <stdio.h>
2 #include <openacc.h>
3 #include <gomp-constants.h>
4
5 #define N (32*32*32+17)
6
7 #pragma acc routine vector
8 void __attribute__ ((noinline)) vector (int ary[N])
9 {
10 #pragma acc loop vector
11 for (unsigned ix = 0; ix < N; ix++)
12 {
13 if (acc_on_device (acc_device_not_host))
14 {
15 int g, w, v;
16
17 g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
18 w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
19 v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
20 ary[ix] = (g << 16) | (w << 8) | v;
21 }
22 else
23 ary[ix] = ix;
24 }
25 }
26
27 int main ()
28 {
29 int ary[N];
30 int ix;
31 int exit = 0;
32 int ondev = 0;
33 int vectorsize;
34
35 for (ix = 0; ix < N;ix++)
36 ary[ix] = -1;
37
38 #define VL 32
39 #pragma acc parallel vector_length(VL) \
40 copy(ary) copy(ondev)
41 {
42 ondev = acc_on_device (acc_device_not_host);
43 vector (ary);
44 }
45 vectorsize = VL;
46 #ifdef ACC_DEVICE_TYPE_radeon
47 /* AMD GCN uses the autovectorizer for the vector dimension: the use
48 of a function call in vector-partitioned code in this test is not
49 currently supported. */
50 vectorsize = 1;
51 #endif
52
53 for (ix = 0; ix < N; ix++)
54 {
55 int expected = ix;
56 if(ondev)
57 {
58 int g = 0;
59 int w = 0;
60 int v = ix % vectorsize;
61
62 expected = (g << 16) | (w << 8) | v;
63 }
64
65 if (ary[ix] != expected)
66 {
67 exit = 1;
68 printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected);
69 }
70 }
71
72 return exit;
73 }