1 /* { dg-additional-options "-fopt-info-note-omp" }
2 { dg-additional-options "--param=openacc-privatization=noisy" }
3 { dg-additional-options "-foffload=-fopt-info-note-omp" }
4 { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
5 for testing/documenting aspects of that functionality. */
6
7 /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
8 aspects of that functionality. */
9
10 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
11 passed to 'incr' may be unset, and in that case, it will be set to [...]",
12 so to maintain compatibility with earlier Tcl releases, we manually
13 initialize counter variables:
14 { dg-line l_dummy[variable c_compute 0 c_loop 0] }
15 { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
16 "WARNING: dg-line var l_dummy defined, but not used". */
17
18 #include <assert.h>
19 #include <openacc.h>
20
21 typedef struct {
22 int x, y;
23 } vec2;
24
25 typedef struct {
26 int x, y, z;
27 int attr[13];
28 } vec3_attr;
29
30
31 /* Test of gang-private variables declared in local scope with parallel
32 directive. */
33
34 void local_g_1()
35 {
36 int i, arr[32];
37
38 for (i = 0; i < 32; i++)
39 arr[i] = 3;
40
41 #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
42 /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
43 /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } l_compute$c_compute } */
44 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
45 {
46 int x;
47
48 #pragma acc loop gang(static:1) /* { dg-line l_loop[incr c_loop] } */
49 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
50 for (i = 0; i < 32; i++)
51 x = i * 2;
52
53 #pragma acc loop gang(static:1) /* { dg-line l_loop[incr c_loop] } */
54 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
55 for (i = 0; i < 32; i++)
56 {
57 if (acc_on_device (acc_device_host))
58 x = i * 2;
59 arr[i] += x;
60 }
61 }
62
63 for (i = 0; i < 32; i++)
64 assert (arr[i] == 3 + i * 2);
65 }
66
67
68 /* Test of worker-private variables declared in a local scope, broadcasting
69 to vector-partitioned mode. Back-to-back worker loops. */
70
71 void local_w_1()
72 {
73 int i, arr[32 * 32 * 32];
74
75 for (i = 0; i < 32 * 32 * 32; i++)
76 arr[i] = i;
77
78 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
79 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
80 {
81 int j;
82
83 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
84 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
85 for (i = 0; i < 32; i++)
86 {
87 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
88 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
89 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
90 /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
91 for (j = 0; j < 32; j++)
92 {
93 int k;
94 int x = i ^ j * 3;
95
96 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
97 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
98 for (k = 0; k < 32; k++)
99 arr[i * 1024 + j * 32 + k] += x * k;
100 }
101
102 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
103 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
104 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
105 /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
106 for (j = 0; j < 32; j++)
107 {
108 int k;
109 int x = i | j * 5;
110
111 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
112 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
113 for (k = 0; k < 32; k++)
114 arr[i * 1024 + j * 32 + k] += x * k;
115 }
116 }
117 }
118
119 for (i = 0; i < 32; i++)
120 for (int j = 0; j < 32; j++)
121 for (int k = 0; k < 32; k++)
122 {
123 int idx = i * 1024 + j * 32 + k;
124 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
125 }
126 }
127
128
129 /* Test of worker-private variables declared in a local scope, broadcasting
130 to vector-partitioned mode. Successive vector loops. */
131
132 void local_w_2()
133 {
134 int i, arr[32 * 32 * 32];
135
136 for (i = 0; i < 32 * 32 * 32; i++)
137 arr[i] = i;
138
139 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
140 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
141 {
142 int j;
143
144 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
145 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
146 for (i = 0; i < 32; i++)
147 {
148 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
149 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
150 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
151 /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
152 for (j = 0; j < 32; j++)
153 {
154 int k;
155 int x = i ^ j * 3;
156
157 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
158 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
159 for (k = 0; k < 32; k++)
160 arr[i * 1024 + j * 32 + k] += x * k;
161
162 x = i | j * 5;
163
164 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
165 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
166 for (k = 0; k < 32; k++)
167 arr[i * 1024 + j * 32 + k] += x * k;
168 }
169 }
170 }
171
172 for (i = 0; i < 32; i++)
173 for (int j = 0; j < 32; j++)
174 for (int k = 0; k < 32; k++)
175 {
176 int idx = i * 1024 + j * 32 + k;
177 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
178 }
179 }
180
181
182 /* Test of worker-private variables declared in a local scope, broadcasting
183 to vector-partitioned mode. Aggregate worker variable. */
184
185 void local_w_3()
186 {
187 int i, arr[32 * 32 * 32];
188
189 for (i = 0; i < 32 * 32 * 32; i++)
190 arr[i] = i;
191
192 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
193 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
194 {
195 int j;
196
197 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
198 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
199 for (i = 0; i < 32; i++)
200 {
201 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
202 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
203 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
204 /* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
205 for (j = 0; j < 32; j++)
206 {
207 int k;
208 vec2 pt;
209
210 pt.x = i ^ j * 3;
211 pt.y = i | j * 5;
212
213 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
214 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
215 for (k = 0; k < 32; k++)
216 arr[i * 1024 + j * 32 + k] += pt.x * k;
217
218 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
219 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
220 for (k = 0; k < 32; k++)
221 arr[i * 1024 + j * 32 + k] += pt.y * k;
222 }
223 }
224 }
225
226 for (i = 0; i < 32; i++)
227 for (int j = 0; j < 32; j++)
228 for (int k = 0; k < 32; k++)
229 {
230 int idx = i * 1024 + j * 32 + k;
231 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
232 }
233 }
234
235
236 /* Test of worker-private variables declared in a local scope, broadcasting
237 to vector-partitioned mode. Addressable worker variable. */
238
239 void local_w_4()
240 {
241 int i, arr[32 * 32 * 32];
242
243 for (i = 0; i < 32 * 32 * 32; i++)
244 arr[i] = i;
245
246 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
247 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
248 {
249 int j;
250
251 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
252 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
253 for (i = 0; i < 32; i++)
254 {
255 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
256 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
257 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
258 /* { dg-note {variable 'pt' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
259 { dg-note {variable 'pt' ought to be adjusted for OpenACC privatization level: 'worker'} "" { target *-*-* } l_loop$c_loop }
260 { dg-note {variable 'pt' adjusted for OpenACC privatization level: 'worker'} "TODO" { target { ! openacc_host_selected } xfail *-*-* } l_loop$c_loop } */
261 /* { dg-note {variable 'ptp' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
262 for (j = 0; j < 32; j++)
263 {
264 int k;
265 vec2 pt, *ptp;
266
267 ptp = &pt;
268
269 pt.x = i ^ j * 3;
270
271 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
272 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
273 for (k = 0; k < 32; k++)
274 arr[i * 1024 + j * 32 + k] += ptp->x * k;
275
276 ptp->y = i | j * 5;
277
278 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
279 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
280 for (k = 0; k < 32; k++)
281 arr[i * 1024 + j * 32 + k] += pt.y * k;
282 }
283 }
284 }
285
286 for (i = 0; i < 32; i++)
287 for (int j = 0; j < 32; j++)
288 for (int k = 0; k < 32; k++)
289 {
290 int idx = i * 1024 + j * 32 + k;
291 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
292 }
293 }
294
295
296 /* Test of worker-private variables declared in a local scope, broadcasting
297 to vector-partitioned mode. Array worker variable. */
298
299 void local_w_5()
300 {
301 int i, arr[32 * 32 * 32];
302
303 for (i = 0; i < 32 * 32 * 32; i++)
304 arr[i] = i;
305
306 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
307 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
308 {
309 int j;
310
311 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
312 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
313 for (i = 0; i < 32; i++)
314 {
315 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
316 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
317 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
318 /* { dg-note {variable 'pt' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
319 for (j = 0; j < 32; j++)
320 {
321 int k;
322 int pt[2];
323
324 pt[0] = i ^ j * 3;
325
326 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
327 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
328 for (k = 0; k < 32; k++)
329 arr[i * 1024 + j * 32 + k] += pt[0] * k;
330
331 pt[1] = i | j * 5;
332
333 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
334 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
335 for (k = 0; k < 32; k++)
336 arr[i * 1024 + j * 32 + k] += pt[1] * k;
337 }
338 }
339 }
340
341 for (i = 0; i < 32; i++)
342 for (int j = 0; j < 32; j++)
343 for (int k = 0; k < 32; k++)
344 {
345 int idx = i * 1024 + j * 32 + k;
346 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
347 }
348 }
349
350
351 /* Test of gang-private variables declared on loop directive. */
352
353 void loop_g_1()
354 {
355 int x = 5, i, arr[32];
356
357 for (i = 0; i < 32; i++)
358 arr[i] = i;
359
360 #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
361 /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } l_compute$c_compute } */
362 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
363 {
364 #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
365 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
366 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
367 for (i = 0; i < 32; i++)
368 {
369 x = i * 2;
370 arr[i] += x;
371 }
372 }
373
374 for (i = 0; i < 32; i++)
375 assert (arr[i] == i * 3);
376 }
377
378
379 /* Test of gang-private variables declared on loop directive, with broadcasting
380 to partitioned workers. */
381
382 void loop_g_2()
383 {
384 int x = 5, i, arr[32 * 32];
385
386 for (i = 0; i < 32 * 32; i++)
387 arr[i] = i;
388
389 #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
390 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
391 {
392 #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
393 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
394 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
395 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
396 for (i = 0; i < 32; i++)
397 {
398 x = i * 2;
399
400 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
401 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
402 for (int j = 0; j < 32; j++)
403 arr[i * 32 + j] += x;
404 }
405 }
406
407 for (i = 0; i < 32 * 32; i++)
408 assert (arr[i] == i + (i / 32) * 2);
409 }
410
411
412 /* Test of gang-private variables declared on loop directive, with broadcasting
413 to partitioned vectors. */
414
415 void loop_g_3()
416 {
417 int x = 5, i, arr[32 * 32];
418
419 for (i = 0; i < 32 * 32; i++)
420 arr[i] = i;
421
422 #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
423 /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } l_compute$c_compute } */
424 {
425 #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
426 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
427 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
428 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
429 for (i = 0; i < 32; i++)
430 {
431 x = i * 2;
432
433 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
434 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
435 for (int j = 0; j < 32; j++)
436 arr[i * 32 + j] += x;
437 }
438 }
439
440 for (i = 0; i < 32 * 32; i++)
441 assert (arr[i] == i + (i / 32) * 2);
442 }
443
444
445 /* Test of gang-private addressable variable declared on loop directive, with
446 broadcasting to partitioned workers. */
447
448 void loop_g_4()
449 {
450 int x = 5, i, arr[32 * 32];
451
452 for (i = 0; i < 32 * 32; i++)
453 arr[i] = i;
454
455 #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
456 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
457 {
458 #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
459 /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
460 But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
461 No longer having address taken: x
462 Now a gimple register: x
463 However, 'x' remains in the candidate set:
464 { dg-note {variable 'x' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
465 Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
466 { dg-note {variable 'x' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_loop$c_loop }
467 For nvptx offloading however, we first mark up 'x', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
468 { dg-note {variable 'x' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_loop$c_loop }
469 { dg-bogus {note: variable 'x' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_loop$c_loop }
470 */
471 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
472 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
473 /* { dg-note {variable 'p' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
474 for (i = 0; i < 32; i++)
475 {
476 int *p = &x;
477
478 x = i * 2;
479
480 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
481 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
482 for (int j = 0; j < 32; j++)
483 arr[i * 32 + j] += x;
484
485 (*p)--;
486 }
487 }
488
489 for (i = 0; i < 32 * 32; i++)
490 assert (arr[i] == i + (i / 32) * 2);
491 }
492
493
494 /* Test of gang-private array variable declared on loop directive, with
495 broadcasting to partitioned workers. */
496
497 void loop_g_5()
498 {
499 int x[8], i, arr[32 * 32];
500
501 for (i = 0; i < 32 * 32; i++)
502 arr[i] = i;
503
504 #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
505 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
506 {
507 #pragma acc loop gang private(x) /* { dg-line l_loop[incr c_loop] } */
508 /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
509 { dg-note {variable 'x' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_loop$c_loop }
510 { dg-note {variable 'x' adjusted for OpenACC privatization level: 'gang'} "" { target { ! openacc_host_selected } } l_loop$c_loop } */
511 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
512 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
513 for (i = 0; i < 32; i++)
514 {
515 for (int j = 0; j < 8; j++)
516 x[j] = j * 2;
517
518 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
519 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
520 for (int j = 0; j < 32; j++)
521 arr[i * 32 + j] += x[j % 8];
522 }
523 }
524
525 for (i = 0; i < 32 * 32; i++)
526 assert (arr[i] == i + (i % 8) * 2);
527 }
528
529
530 /* Test of gang-private aggregate variable declared on loop directive, with
531 broadcasting to partitioned workers. */
532
533 void loop_g_6()
534 {
535 int i, arr[32 * 32];
536 vec3_attr pt;
537
538 for (i = 0; i < 32 * 32; i++)
539 arr[i] = i;
540
541 #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
542 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
543 {
544 #pragma acc loop gang private(pt) /* { dg-line l_loop[incr c_loop] } */
545 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
546 /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
547 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
548 for (i = 0; i < 32; i++)
549 {
550 pt.x = i;
551 pt.y = i * 2;
552 pt.z = i * 4;
553 pt.attr[5] = i * 6;
554
555 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
556 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
557 for (int j = 0; j < 32; j++)
558 arr[i * 32 + j] += pt.x + pt.y + pt.z + pt.attr[5];
559 }
560 }
561
562 for (i = 0; i < 32 * 32; i++)
563 assert (arr[i] == i + (i / 32) * 13);
564 }
565
566
567 /* Test of vector-private variables declared on loop directive. */
568
569 void loop_v_1()
570 {
571 int x, i, arr[32 * 32 * 32];
572
573 for (i = 0; i < 32 * 32 * 32; i++)
574 arr[i] = i;
575
576 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
577 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
578 {
579 int j;
580
581 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
582 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
583 for (i = 0; i < 32; i++)
584 {
585 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
586 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
587 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
588 for (j = 0; j < 32; j++)
589 {
590 int k;
591
592 #pragma acc loop vector private(x) /* { dg-line l_loop[incr c_loop] } */
593 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
594 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
595 for (k = 0; k < 32; k++)
596 {
597 x = i ^ j * 3;
598 arr[i * 1024 + j * 32 + k] += x * k;
599 }
600
601 #pragma acc loop vector private(x) /* { dg-line l_loop[incr c_loop] } */
602 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
603 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
604 for (k = 0; k < 32; k++)
605 {
606 x = i | j * 5;
607 arr[i * 1024 + j * 32 + k] += x * k;
608 }
609 }
610 }
611 }
612
613 for (i = 0; i < 32; i++)
614 for (int j = 0; j < 32; j++)
615 for (int k = 0; k < 32; k++)
616 {
617 int idx = i * 1024 + j * 32 + k;
618 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
619 }
620 }
621
622
623 /* Test of vector-private variables declared on loop directive. Array type. */
624
625 void loop_v_2()
626 {
627 int pt[2], i, arr[32 * 32 * 32];
628
629 for (i = 0; i < 32 * 32 * 32; i++)
630 arr[i] = i;
631
632 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
633 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
634 {
635 int j;
636
637 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
638 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
639 for (i = 0; i < 32; i++)
640 {
641 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
642 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
643 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
644 for (j = 0; j < 32; j++)
645 {
646 int k;
647
648 #pragma acc loop vector private(pt) /* { dg-line l_loop[incr c_loop] } */
649 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
650 /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
651 for (k = 0; k < 32; k++)
652 {
653 pt[0] = i ^ j * 3;
654 pt[1] = i | j * 5;
655 arr[i * 1024 + j * 32 + k] += pt[0] * k;
656 arr[i * 1024 + j * 32 + k] += pt[1] * k;
657 }
658 }
659 }
660 }
661
662 for (i = 0; i < 32; i++)
663 for (int j = 0; j < 32; j++)
664 for (int k = 0; k < 32; k++)
665 {
666 int idx = i * 1024 + j * 32 + k;
667 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
668 }
669 }
670
671
672 /* Test of worker-private variables declared on a loop directive. */
673
674 void loop_w_1()
675 {
676 int x = 5, i, arr[32 * 32];
677
678 for (i = 0; i < 32 * 32; i++)
679 arr[i] = i;
680
681 #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
682 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
683 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
684 {
685 int j;
686
687 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
688 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
689 for (i = 0; i < 32; i++)
690 {
691 #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
692 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
693 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
694 for (j = 0; j < 32; j++)
695 {
696 x = i ^ j * 3;
697 /* Try to ensure 'x' accesses doesn't get optimized into a
698 temporary. */
699 __asm__ __volatile__ ("");
700 arr[i * 32 + j] += x;
701 }
702 }
703 }
704
705 for (i = 0; i < 32 * 32; i++)
706 assert (arr[i] == i + ((i / 32) ^ (i % 32) * 3));
707 }
708
709
710 /* Test of worker-private variables declared on a loop directive, broadcasting
711 to vector-partitioned mode. */
712
713 void loop_w_2()
714 {
715 int x = 5, i, arr[32 * 32 * 32];
716
717 for (i = 0; i < 32 * 32 * 32; i++)
718 arr[i] = i;
719
720 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
721 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
722 {
723 int j;
724
725 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
726 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
727 for (i = 0; i < 32; i++)
728 {
729 #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
730 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
731 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
732 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
733 for (j = 0; j < 32; j++)
734 {
735 int k;
736 x = i ^ j * 3;
737
738 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
739 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
740 for (k = 0; k < 32; k++)
741 arr[i * 1024 + j * 32 + k] += x * k;
742 }
743 }
744 }
745
746 for (i = 0; i < 32; i++)
747 for (int j = 0; j < 32; j++)
748 for (int k = 0; k < 32; k++)
749 {
750 int idx = i * 1024 + j * 32 + k;
751 assert (arr[idx] == idx + (i ^ j * 3) * k);
752 }
753 }
754
755
756 /* Test of worker-private variables declared on a loop directive, broadcasting
757 to vector-partitioned mode. Back-to-back worker loops. */
758
759 void loop_w_3()
760 {
761 int x = 5, i, arr[32 * 32 * 32];
762
763 for (i = 0; i < 32 * 32 * 32; i++)
764 arr[i] = i;
765
766 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
767 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
768 {
769 int j;
770
771 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
772 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
773 for (i = 0; i < 32; i++)
774 {
775 #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
776 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
777 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
778 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
779 for (j = 0; j < 32; j++)
780 {
781 int k;
782 x = i ^ j * 3;
783
784 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
785 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
786 for (k = 0; k < 32; k++)
787 arr[i * 1024 + j * 32 + k] += x * k;
788 }
789
790 #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
791 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
792 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
793 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
794 for (j = 0; j < 32; j++)
795 {
796 int k;
797 x = i | j * 5;
798
799 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
800 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
801 for (k = 0; k < 32; k++)
802 arr[i * 1024 + j * 32 + k] += x * k;
803 }
804 }
805 }
806
807 for (i = 0; i < 32; i++)
808 for (int j = 0; j < 32; j++)
809 for (int k = 0; k < 32; k++)
810 {
811 int idx = i * 1024 + j * 32 + k;
812 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
813 }
814 }
815
816
817 /* Test of worker-private variables declared on a loop directive, broadcasting
818 to vector-partitioned mode. Successive vector loops. */
819
820 void loop_w_4()
821 {
822 int x = 5, i, arr[32 * 32 * 32];
823
824 for (i = 0; i < 32 * 32 * 32; i++)
825 arr[i] = i;
826
827 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
828 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
829 {
830 int j;
831
832 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
833 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
834 for (i = 0; i < 32; i++)
835 {
836 #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
837 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
838 /* { dg-note {variable 'x' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
839 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
840 for (j = 0; j < 32; j++)
841 {
842 int k;
843 x = i ^ j * 3;
844
845 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
846 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
847 for (k = 0; k < 32; k++)
848 arr[i * 1024 + j * 32 + k] += x * k;
849
850 x = i | j * 5;
851
852 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
853 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
854 for (k = 0; k < 32; k++)
855 arr[i * 1024 + j * 32 + k] += x * k;
856 }
857 }
858 }
859
860 for (i = 0; i < 32; i++)
861 for (int j = 0; j < 32; j++)
862 for (int k = 0; k < 32; k++)
863 {
864 int idx = i * 1024 + j * 32 + k;
865 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
866 }
867 }
868
869
870 /* Test of worker-private variables declared on a loop directive, broadcasting
871 to vector-partitioned mode. Addressable worker variable. */
872
873 void loop_w_5()
874 {
875 int x = 5, i, arr[32 * 32 * 32];
876
877 for (i = 0; i < 32 * 32 * 32; i++)
878 arr[i] = i;
879
880 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
881 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
882 {
883 int j;
884
885 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
886 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
887 for (i = 0; i < 32; i++)
888 {
889 #pragma acc loop worker private(x) /* { dg-line l_loop[incr c_loop] } */
890 /* { dg-note {variable 'x' in 'private' clause is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
891 { dg-note {variable 'x' ought to be adjusted for OpenACC privatization level: 'worker'} "" { target *-*-* } l_loop$c_loop }
892 { dg-note {variable 'x' adjusted for OpenACC privatization level: 'worker'} "TODO" { target { ! openacc_host_selected } xfail *-*-* } l_loop$c_loop } */
893 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
894 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
895 /* { dg-note {variable 'p' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
896 for (j = 0; j < 32; j++)
897 {
898 int k;
899 int *p = &x;
900
901 x = i ^ j * 3;
902
903 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
904 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
905 for (k = 0; k < 32; k++)
906 arr[i * 1024 + j * 32 + k] += x * k;
907
908 *p = i | j * 5;
909
910 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
911 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
912 for (k = 0; k < 32; k++)
913 arr[i * 1024 + j * 32 + k] += x * k;
914 }
915 }
916 }
917
918 for (i = 0; i < 32; i++)
919 for (int j = 0; j < 32; j++)
920 for (int k = 0; k < 32; k++)
921 {
922 int idx = i * 1024 + j * 32 + k;
923 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
924 }
925 }
926
927
928 /* Test of worker-private variables declared on a loop directive, broadcasting
929 to vector-partitioned mode. Aggregate worker variable. */
930
931 void loop_w_6()
932 {
933 int i, arr[32 * 32 * 32];
934 vec2 pt;
935
936 for (i = 0; i < 32 * 32 * 32; i++)
937 arr[i] = i;
938
939 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
940 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
941 {
942 int j;
943
944 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
945 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
946 for (i = 0; i < 32; i++)
947 {
948 #pragma acc loop worker private(pt) /* { dg-line l_loop[incr c_loop] } */
949 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
950 /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
951 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
952 for (j = 0; j < 32; j++)
953 {
954 int k;
955
956 pt.x = i ^ j * 3;
957 pt.y = i | j * 5;
958
959 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
960 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
961 for (k = 0; k < 32; k++)
962 arr[i * 1024 + j * 32 + k] += pt.x * k;
963
964 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
965 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
966 for (k = 0; k < 32; k++)
967 arr[i * 1024 + j * 32 + k] += pt.y * k;
968 }
969 }
970 }
971
972 for (i = 0; i < 32; i++)
973 for (int j = 0; j < 32; j++)
974 for (int k = 0; k < 32; k++)
975 {
976 int idx = i * 1024 + j * 32 + k;
977 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
978 }
979 }
980
981
982 /* Test of worker-private variables declared on loop directive, broadcasting
983 to vector-partitioned mode. Array worker variable. */
984
985 void loop_w_7()
986 {
987 int i, arr[32 * 32 * 32];
988 int pt[2];
989
990 for (i = 0; i < 32 * 32 * 32; i++)
991 arr[i] = i;
992
993 /* "pt" is treated as "present_or_copy" on the parallel directive because it
994 is an array variable. */
995 #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
996 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
997 {
998 int j;
999
1000 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
1001 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1002 for (i = 0; i < 32; i++)
1003 {
1004 /* But here, it is made private per-worker. */
1005 #pragma acc loop worker private(pt) /* { dg-line l_loop[incr c_loop] } */
1006 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1007 /* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1008 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1009 for (j = 0; j < 32; j++)
1010 {
1011 int k;
1012
1013 pt[0] = i ^ j * 3;
1014
1015 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
1016 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1017 for (k = 0; k < 32; k++)
1018 arr[i * 1024 + j * 32 + k] += pt[0] * k;
1019
1020 pt[1] = i | j * 5;
1021
1022 #pragma acc loop vector /* { dg-line l_loop[incr c_loop] } */
1023 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1024 for (k = 0; k < 32; k++)
1025 arr[i * 1024 + j * 32 + k] += pt[1] * k;
1026 }
1027 }
1028 }
1029
1030 for (i = 0; i < 32; i++)
1031 for (int j = 0; j < 32; j++)
1032 for (int k = 0; k < 32; k++)
1033 {
1034 int idx = i * 1024 + j * 32 + k;
1035 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
1036 }
1037 }
1038
1039
1040 /* Test of gang-private variables declared on the parallel directive. */
1041
1042 void parallel_g_1()
1043 {
1044 int x = 5, i, arr[32];
1045
1046 for (i = 0; i < 32; i++)
1047 arr[i] = 3;
1048
1049 #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
1050 /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } l_compute$c_compute } */
1051 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
1052 {
1053 #pragma acc loop gang(static:1) /* { dg-line l_loop[incr c_loop] } */
1054 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1055 for (i = 0; i < 32; i++)
1056 x = i * 2;
1057
1058 #pragma acc loop gang(static:1) /* { dg-line l_loop[incr c_loop] } */
1059 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1060 for (i = 0; i < 32; i++)
1061 {
1062 if (acc_on_device (acc_device_host))
1063 x = i * 2;
1064 arr[i] += x;
1065 }
1066 }
1067
1068 for (i = 0; i < 32; i++)
1069 assert (arr[i] == 3 + i * 2);
1070 }
1071
1072
1073 /* Test of gang-private array variable declared on the parallel directive. */
1074
1075 void parallel_g_2()
1076 {
1077 int x[32], i, arr[32 * 32];
1078
1079 for (i = 0; i < 32 * 32; i++)
1080 arr[i] = i;
1081
1082 #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(2) vector_length(32) /* { dg-line l_compute[incr c_compute] } */
1083 /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } l_compute$c_compute } */
1084 {
1085 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
1086 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1087 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1088 for (i = 0; i < 32; i++)
1089 {
1090 int j;
1091 for (j = 0; j < 32; j++)
1092 x[j] = j * 2;
1093
1094 #pragma acc loop worker /* { dg-line l_loop[incr c_loop] } */
1095 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
1096 for (j = 0; j < 32; j++)
1097 arr[i * 32 + j] += x[31 - j];
1098 }
1099 }
1100
1101 for (i = 0; i < 32 * 32; i++)
1102 assert (arr[i] == i + (31 - (i % 32)) * 2);
1103 }
1104
1105
1106 int main ()
1107 {
1108 local_g_1();
1109 local_w_1();
1110 local_w_2();
1111 local_w_3();
1112 local_w_4();
1113 local_w_5();
1114 loop_g_1();
1115 loop_g_2();
1116 loop_g_3();
1117 loop_g_4();
1118 loop_g_5();
1119 loop_g_6();
1120 loop_v_1();
1121 loop_v_2();
1122 loop_w_1();
1123 loop_w_2();
1124 loop_w_3();
1125 loop_w_4();
1126 loop_w_5();
1127 loop_w_6();
1128 loop_w_7();
1129 parallel_g_1();
1130 parallel_g_2();
1131
1132 return 0;
1133 }