1 /* { dg-additional-options "-fopenacc-dim=32" } */
2
3 #include <stdio.h>
4 #include <openacc.h>
5 #include <gomp-constants.h>
6
check(const int * ary,int size,int gp,int wp,int vp)7 static int check (const int *ary, int size, int gp, int wp, int vp)
8 {
9 int exit = 0;
10 int ix;
11 int gangs[32], workers[32], vectors[32];
12
13 for (ix = 0; ix < 32; ix++)
14 gangs[ix] = workers[ix] = vectors[ix] = 0;
15
16 for (ix = 0; ix < size; ix++)
17 {
18 vectors[ary[ix] & 0xff]++;
19 workers[(ary[ix] >> 8) & 0xff]++;
20 gangs[(ary[ix] >> 16) & 0xff]++;
21 }
22
23 for (ix = 0; ix < 32; ix++)
24 {
25 if (gp)
26 {
27 int expect = gangs[0];
28 if (gangs[ix] != expect)
29 {
30 exit = 1;
31 printf ("gang %d not used %d times\n", ix, expect);
32 }
33 }
34 else if (ix && gangs[ix])
35 {
36 exit = 1;
37 printf ("gang %d unexpectedly used\n", ix);
38 }
39
40 if (wp)
41 {
42 int expect = workers[0];
43 if (workers[ix] != expect)
44 {
45 exit = 1;
46 printf ("worker %d not used %d times\n", ix, expect);
47 }
48 }
49 else if (ix && workers[ix])
50 {
51 exit = 1;
52 printf ("worker %d unexpectedly used\n", ix);
53 }
54
55 if (vp)
56 {
57 int expect = vectors[0];
58 if (vectors[ix] != expect)
59 {
60 exit = 1;
61 printf ("vector %d not used %d times\n", ix, expect);
62 }
63 }
64 else if (ix && vectors[ix])
65 {
66 exit = 1;
67 printf ("vector %d unexpectedly used\n", ix);
68 }
69
70 }
71 return exit;
72 }
73
74 #pragma acc routine seq
place()75 static int __attribute__((noinline)) place ()
76 {
77 int r = 0;
78
79 if (acc_on_device (acc_device_not_host))
80 {
81 int g, w, v;
82
83 g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
84 w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
85 v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
86 r = (g << 16) | (w << 8) | v;
87 }
88 return r;
89 }
90
clear(int * ary,int size)91 static void clear (int *ary, int size)
92 {
93 int ix;
94
95 for (ix = 0; ix < size; ix++)
96 ary[ix] = -1;
97 }
98
gang_vector_1(int * ary,int size)99 int gang_vector_1 (int *ary, int size)
100 {
101 clear (ary, size);
102 #pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
103 {
104 #pragma acc loop tile(128) gang vector
105 for (int jx = 0; jx < size; jx++)
106 ary[jx] = place ();
107 }
108
109 return check (ary, size, 1, 0, 1);
110 }
111
gang_vector_2a(int * ary,int size)112 int gang_vector_2a (int *ary, int size)
113 {
114 if (size % 256)
115 return 1;
116
117 clear (ary, size);
118 #pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
119 {
120 #pragma acc loop tile(64, 64) gang vector
121 for (int jx = 0; jx < size / 256; jx++)
122 for (int ix = 0; ix < 256; ix++)
123 ary[jx * 256 + ix] = place ();
124 }
125
126 return check (ary, size, 1, 0, 1);
127 }
128
gang_vector_2b(int * ary,int size)129 int gang_vector_2b (int *ary, int size)
130 {
131 if (size % 256)
132 return 1;
133
134 clear (ary, size);
135 #pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size)
136 {
137 #pragma acc loop tile(64, 64) gang vector
138 for (int jx = 0; jx < size; jx += 256)
139 for (int ix = 0; ix < 256; ix++)
140 ary[jx + ix] = place ();
141 }
142
143 return check (ary, size, 1, 0, 1);
144 }
145
worker_vector_2a(int * ary,int size)146 int worker_vector_2a (int *ary, int size)
147 {
148 if (size % 256)
149 return 1;
150
151 clear (ary, size);
152 #pragma acc parallel vector_length(32) num_workers (32) copy (ary[0:size]) firstprivate (size)
153 {
154 #pragma acc loop tile(64, 64) worker vector
155 for (int jx = 0; jx < size / 256; jx++)
156 for (int ix = 0; ix < 256; ix++)
157 ary[jx * 256 + ix] = place ();
158 }
159
160 return check (ary, size, 0, 1, 1);
161 }
162
worker_vector_2b(int * ary,int size)163 int worker_vector_2b (int *ary, int size)
164 {
165 if (size % 256)
166 return 1;
167
168 clear (ary, size);
169 #pragma acc parallel vector_length(32) num_workers (32) copy (ary[0:size]) firstprivate (size)
170 {
171 #pragma acc loop tile(64, 64) worker vector
172 for (int jx = 0; jx < size; jx += 256)
173 for (int ix = 0; ix < 256; ix++)
174 ary[jx + ix] = place ();
175 }
176
177 return check (ary, size, 0, 1, 1);
178 }
179
gang_worker_vector_2a(int * ary,int size)180 int gang_worker_vector_2a (int *ary, int size)
181 {
182 if (size % 256)
183 return 1;
184 clear (ary, size);
185 #pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
186 {
187 #pragma acc loop tile(32, 32)
188 for (int jx = 0; jx < size / 256; jx++)
189 for (int ix = 0; ix < 256; ix++)
190 ary[jx * 256 + ix] = place ();
191 }
192
193 return check (ary, size, 1, 1, 1);
194 }
195
gang_worker_vector_2b(int * ary,int size)196 int gang_worker_vector_2b (int *ary, int size)
197 {
198 if (size % 256)
199 return 1;
200 clear (ary, size);
201 #pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
202 {
203 #pragma acc loop tile(32, 32)
204 for (int jx = 0; jx < size; jx += 256)
205 for (int ix = 0; ix < 256; ix++)
206 ary[jx + ix] = place ();
207 }
208
209 return check (ary, size, 1, 1, 1);
210 }
211
gang_worker_vector_star_2a(int * ary,int size)212 int gang_worker_vector_star_2a (int *ary, int size)
213 {
214 if (size % 256)
215 return 1;
216
217 clear (ary, size);
218 #pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
219 {
220 #pragma acc loop tile(*, *)
221 for (int jx = 0; jx < size / 256; jx++)
222 for (int ix = 0; ix < 256; ix++)
223 ary[jx * 256 + ix] = place ();
224 }
225
226 return check (ary, size, 1, 1, 1);
227 }
228
gang_worker_vector_star_2b(int * ary,int size)229 int gang_worker_vector_star_2b (int *ary, int size)
230 {
231 if (size % 256)
232 return 1;
233
234 clear (ary, size);
235 #pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size)
236 {
237 #pragma acc loop tile(*, *)
238 for (int jx = 0; jx < size; jx +=256)
239 for (int ix = 0; ix < 256; ix++)
240 ary[jx + ix] = place ();
241 }
242
243 return check (ary, size, 1, 1, 1);
244 }
245
246 #define N (32*32*32*8)
main()247 int main ()
248 {
249 int ondev = 0;
250
251 #pragma acc parallel copy(ondev)
252 {
253 ondev = acc_on_device (acc_device_not_host);
254 }
255 if (!ondev)
256 return 0;
257
258 int ary[N];
259 if (gang_vector_1 (ary, N))
260 return 1;
261 if (gang_vector_2a (ary, N))
262 return 1;
263 if (worker_vector_2a (ary, N))
264 return 1;
265 if (gang_worker_vector_2a (ary, N))
266 return 1;
267 if (gang_worker_vector_star_2a (ary, N))
268 return 1;
269 if (gang_vector_2b (ary, N))
270 return 1;
271 if (worker_vector_2b (ary, N))
272 return 1;
273 if (gang_worker_vector_2b (ary, N))
274 return 1;
275 if (gang_worker_vector_star_2b (ary, N))
276 return 1;
277 return 0;
278 }
279