]>
Commit | Line | Data |
---|---|---|
2620c80d TS |
1 | #include <stdio.h> |
2 | #include <openacc.h> | |
1f62d637 | 3 | #include <gomp-constants.h> |
2620c80d TS |
4 | |
5 | #define NUM_WORKERS 16 | |
6 | #define NUM_VECTORS 32 | |
7 | #define WIDTH 64 | |
8 | #define HEIGHT 32 | |
9 | ||
10 | #define WORK_ID(I,N) \ | |
1f62d637 TV |
11 | (acc_on_device (acc_device_not_host) \ |
12 | ? __builtin_goacc_parlevel_id (GOMP_DIM_WORKER) \ | |
13 | : (I % N)) | |
2620c80d | 14 | #define VEC_ID(I,N) \ |
1f62d637 TV |
15 | (acc_on_device (acc_device_not_host) \ |
16 | ? __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR) \ | |
17 | : (I % N)) | |
2620c80d TS |
18 | |
19 | #pragma acc routine worker | |
20 | void __attribute__ ((noinline)) | |
21 | WorkVec (int *ptr, int w, int h, int nw, int nv) | |
22 | { | |
23 | #pragma acc loop worker | |
24 | for (int i = 0; i < h; i++) | |
25 | #pragma acc loop vector | |
26 | for (int j = 0; j < w; j++) | |
27 | ptr[i*w + j] = (WORK_ID (i, nw) << 8) | VEC_ID(j, nv); | |
28 | } | |
29 | ||
30 | int DoWorkVec (int nw) | |
31 | { | |
32 | int ary[HEIGHT][WIDTH]; | |
33 | int err = 0; | |
34 | ||
35 | for (int ix = 0; ix != HEIGHT; ix++) | |
36 | for (int jx = 0; jx != WIDTH; jx++) | |
37 | ary[ix][jx] = 0xdeadbeef; | |
38 | ||
39 | printf ("spawning %d ...", nw); fflush (stdout); | |
40 | ||
41 | #pragma acc parallel num_workers(nw) vector_length (NUM_VECTORS) copy (ary) | |
42 | { | |
43 | WorkVec ((int *)ary, WIDTH, HEIGHT, nw, NUM_VECTORS); | |
44 | } | |
45 | ||
46 | for (int ix = 0; ix != HEIGHT; ix++) | |
47 | for (int jx = 0; jx != WIDTH; jx++) | |
48 | { | |
49 | int exp = ((ix % nw) << 8) | (jx % NUM_VECTORS); | |
50 | ||
51 | if (ary[ix][jx] != exp) | |
52 | { | |
53 | printf ("\nary[%d][%d] = %#x expected %#x", ix, jx, | |
54 | ary[ix][jx], exp); | |
55 | err = 1; | |
56 | } | |
57 | } | |
58 | printf (err ? " failed\n" : " ok\n"); | |
59 | ||
60 | return err; | |
61 | } | |
62 | ||
63 | int main () | |
64 | { | |
65 | int err = 0; | |
66 | ||
67 | for (int W = 1; W <= NUM_WORKERS; W <<= 1) | |
68 | err |= DoWorkVec (W); | |
69 | ||
70 | return err; | |
71 | } |