]>
Commit | Line | Data |
---|---|---|
5d6823a2 | 1 | #include <stdio.h> |
1f62d637 TV |
2 | #include <openacc.h> |
3 | #include <gomp-constants.h> | |
5d6823a2 NS |
4 | |
5 | #define N (32*32*32+17) | |
6 | ||
7 | int main () | |
8 | { | |
9 | int ix; | |
10 | int ondev = 0; | |
11 | int t = 0, h = 0; | |
12 | ||
2c71d454 | 13 | #pragma acc parallel vector_length(32) copy(ondev) |
5d6823a2 NS |
14 | { |
15 | #pragma acc loop vector reduction (+:t) | |
16 | for (unsigned ix = 0; ix < N; ix++) | |
17 | { | |
18 | int val = ix; | |
19 | ||
1f62d637 | 20 | if (acc_on_device (acc_device_not_host)) |
5d6823a2 | 21 | { |
1f62d637 | 22 | int g, w, v; |
5d6823a2 | 23 | |
1f62d637 TV |
24 | g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG); |
25 | w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); | |
26 | v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); | |
5d6823a2 NS |
27 | val = (g << 16) | (w << 8) | v; |
28 | ondev = 1; | |
29 | } | |
30 | t += val; | |
31 | } | |
32 | } | |
33 | ||
34 | for (ix = 0; ix < N; ix++) | |
35 | { | |
36 | int val = ix; | |
37 | if (ondev) | |
38 | { | |
39 | int g = 0; | |
40 | int w = 0; | |
41 | int v = ix % 32; | |
42 | ||
43 | val = (g << 16) | (w << 8) | v; | |
44 | } | |
45 | h += val; | |
46 | } | |
47 | ||
48 | if (t != h) | |
49 | { | |
50 | printf ("t=%x expected %x\n", t, h); | |
51 | return 1; | |
52 | } | |
53 | ||
54 | return 0; | |
55 | } |