]>
Commit | Line | Data |
---|---|---|
36fa0f5f | 1 | #include <stdio.h> |
1f62d637 TV |
2 | #include <openacc.h> |
3 | #include <gomp-constants.h> | |
36fa0f5f NS |
4 | |
5 | #define N (32*32*32+17) | |
6 | ||
7 | #pragma acc routine worker | |
8 | void __attribute__ ((noinline)) worker (int ary[N]) | |
9 | { | |
10 | #pragma acc loop worker vector | |
11 | for (unsigned ix = 0; ix < N; ix++) | |
12 | { | |
1f62d637 | 13 | if (acc_on_device (acc_device_not_host)) |
36fa0f5f | 14 | { |
1f62d637 | 15 | int g, w, v; |
36fa0f5f | 16 | |
1f62d637 TV |
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); | |
36fa0f5f NS |
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 | ||
34 | for (ix = 0; ix < N;ix++) | |
35 | ary[ix] = -1; | |
36 | ||
37 | #pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) | |
38 | { | |
1f62d637 | 39 | ondev = acc_on_device (acc_device_not_host); |
36fa0f5f NS |
40 | worker (ary); |
41 | } | |
42 | ||
43 | for (ix = 0; ix < N; ix++) | |
44 | { | |
45 | int expected = ix; | |
46 | if(ondev) | |
47 | { | |
48 | int g = 0; | |
49 | int w = (ix / 32) % 32; | |
50 | int v = ix % 32; | |
51 | ||
52 | expected = (g << 16) | (w << 8) | v; | |
53 | } | |
54 | ||
55 | if (ary[ix] != expected) | |
56 | { | |
57 | exit = 1; | |
58 | printf ("ary[%d]=%x expected %x\n", ix, ary[ix], expected); | |
59 | } | |
60 | } | |
61 | ||
62 | return exit; | |
63 | } |