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