1 /* { dg-do run } */
2
3 /* Based on '../libgomp.oacc-fortran/asyncwait-1.f90'. */
4
5 /* { dg-additional-options "--param=openacc-kernels=decompose" } */
6
7 /* { dg-additional-options "-fopt-info-all-omp" }
8 { dg-additional-options "-foffload=-fopt-info-all-omp" } */
9
10 /* { dg-additional-options "--param=openacc-privatization=noisy" }
11 { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
12 Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
13 { dg-prune-output {note: variable '[Di]\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
14
15 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
16 passed to 'incr' may be unset, and in that case, it will be set to [...]",
17 so to maintain compatibility with earlier Tcl releases, we manually
18 initialize counter variables:
19 { dg-line l_dummy[variable c_compute 0 c_loop_i 0] }
20 { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
21 "WARNING: dg-line var l_dummy defined, but not used". */
22
23 #include <stdlib.h>
24
25 #define N 64
26
27 int
28 main (void)
29 {
30 int *a, *b, *c, *d, *e;
31
32 a = (int*)malloc (N * sizeof (*a));
33 b = (int*)malloc (N * sizeof (*b));
34 c = (int*)malloc (N * sizeof (*c));
35 d = (int*)malloc (N * sizeof (*d));
36 e = (int*)malloc (N * sizeof (*e));
37
38 for (int i = 0; i < N; ++i)
39 {
40 a[i] = 3;
41 b[i] = 0;
42 }
43
44 #pragma acc data copy (a[0:N]) copy (b[0:N])
45 {
46
47 #pragma acc parallel async /* { dg-line l_compute[incr c_compute] } */
48 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
49 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
50 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
51 /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
52 for (int i = 0; i < N; ++i)
53 b[i] = a[i];
54
55 #pragma acc wait
56 }
57
58 for (int i = 0; i < N; ++i)
59 {
60 if (a[i] != 3)
61 abort ();
62 if (b[i] != 3)
63 abort ();
64 }
65
66 for (int i = 0; i < N; ++i)
67 {
68 a[i] = 2;
69 b[i] = 0;
70 }
71
72 #pragma acc data copy (a[0:N]) copy (b[0:N])
73 {
74 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
75 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
76 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
77 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
78 /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
79 for (int i = 0; i < N; ++i)
80 b[i] = a[i];
81
82 #pragma acc wait (1)
83 }
84
85 for (int i = 0; i < N; ++i)
86 {
87 if (a[i] != 2) abort ();
88 if (b[i] != 2) abort ();
89 }
90
91 for (int i = 0; i < N; ++i)
92 {
93 a[i] = 3;
94 b[i] = 0;
95 c[i] = 0;
96 d[i] = 0;
97 }
98
99 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N])
100 {
101
102 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
103 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
104 for (int i = 0; i < N; ++i)
105 b[i] = (a[i] * a[i] * a[i]) / a[i];
106
107 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
108 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
109 for (int i = 0; i < N; ++i)
110 c[i] = (a[i] * 4) / a[i];
111
112
113 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
114 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
115 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
116 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
117 /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
118 for (int i = 0; i < N; ++i)
119 d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
120
121 #pragma acc wait (1)
122 }
123
124 for (int i = 0; i < N; ++i)
125 {
126 if (a[i] != 3)
127 abort ();
128 if (b[i] != 9)
129 abort ();
130 if (c[i] != 4)
131 abort ();
132 if (d[i] != 1)
133 abort ();
134 }
135
136 for (int i = 0; i < N; ++i)
137 {
138 a[i] = 2;
139 b[i] = 0;
140 c[i] = 0;
141 d[i] = 0;
142 e[i] = 0;
143 }
144
145 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N])
146 {
147
148 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
149 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
150 for (int i = 0; i < N; ++i)
151 b[i] = (a[i] * a[i] * a[i]) / a[i];
152
153 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
154 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
155 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
156 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
157 /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
158 for (int i = 0; i < N; ++i)
159 c[i] = (a[i] * 4) / a[i];
160
161 #pragma acc parallel async (1) /* { dg-line l_compute[incr c_compute] } */
162 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
163 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
164 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
165 /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
166 for (int i = 0; i < N; ++i)
167 d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
168
169
170 #pragma acc parallel wait (1) async (1) /* { dg-line l_compute[incr c_compute] } */
171 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
172 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
173 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
174 /* { dg-optimized "assigned OpenACC gang vector loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
175 for (int i = 0; i < N; ++i)
176 e[i] = a[i] + b[i] + c[i] + d[i];
177
178 #pragma acc wait (1)
179 }
180
181 for (int i = 0; i < N; ++i)
182 {
183 if (a[i] != 2)
184 abort ();
185 if (b[i] != 4)
186 abort ();
187 if (c[i] != 4)
188 abort ();
189 if (d[i] != 1)
190 abort ();
191 if (e[i] != 11)
192 abort ();
193 }
194
195 for (int i = 0; i < N; ++i)
196 {
197 a[i] = 3;
198 b[i] = 0;
199 }
200
201 #pragma acc data copy (a[0:N]) copy (b[0:N])
202 {
203
204 #pragma acc kernels async
205 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
206 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
207 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
208 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
209 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
210 for (int i = 0; i < N; ++i)
211 b[i] = a[i];
212
213 #pragma acc wait
214 }
215
216 for (int i = 0; i < N; ++i)
217 {
218 if (a[i] != 3)
219 abort ();
220 if (b[i] != 3)
221 abort ();
222 }
223
224 for (int i = 0; i < N; ++i)
225 {
226 a[i] = 2;
227 b[i] = 0;
228 }
229
230 #pragma acc data copy (a[0:N]) copy (b[0:N])
231 {
232 #pragma acc kernels async (1)
233 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
234 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
235 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
236 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
237 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
238 for (int i = 0; i < N; ++i)
239 b[i] = a[i];
240
241 #pragma acc wait (1)
242 }
243
244 for (int i = 0; i < N; ++i)
245 {
246 if (a[i] != 2)
247 abort ();
248 if (b[i] != 2)
249 abort ();
250 }
251
252 for (int i = 0; i < N; ++i)
253 {
254 a[i] = 3;
255 b[i] = 0;
256 c[i] = 0;
257 d[i] = 0;
258 }
259
260 #pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N])
261 {
262 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
263 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */
264 /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
265 /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */
266 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute }
267 { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */
268 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
269 for (int i = 0; i < N; ++i)
270 b[i] = (a[i] * a[i] * a[i]) / a[i];
271
272 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
273 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */
274 /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
275 /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */
276 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute }
277 { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */
278 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
279 for (int i = 0; i < N; ++i)
280 c[i] = (a[i] * 4) / a[i];
281
282 #pragma acc kernels async (1)
283 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
284 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
285 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
286 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
287 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
288 for (int i = 0; i < N; ++i)
289 d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
290
291 #pragma acc wait (1)
292 }
293
294 for (int i = 0; i < N; ++i)
295 {
296 if (a[i] != 3)
297 abort ();
298 if (b[i] != 9)
299 abort ();
300 if (c[i] != 4)
301 abort ();
302 if (d[i] != 1)
303 abort ();
304 }
305
306 for (int i = 0; i < N; ++i)
307 {
308 a[i] = 2;
309 b[i] = 0;
310 c[i] = 0;
311 d[i] = 0;
312 e[i] = 0;
313 }
314
315 #pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N])
316 {
317 #pragma acc kernels async (1) /* { dg-line l_compute[incr c_compute] } */
318 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} "" { target *-*-* } l_compute$c_compute } */
319 /* { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
320 /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute } */
321 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target { ! __OPTIMIZE__ } } l_compute$c_compute }
322 { dg-optimized "assigned OpenACC gang loop parallelism" "" { target { __OPTIMIZE__ } } l_compute$c_compute } */
323 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
324 for (int i = 0; i < N; ++i)
325 b[i] = (a[i] * a[i] * a[i]) / a[i];
326
327 #pragma acc kernels async (1)
328 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
329 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
330 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
331 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
332 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
333 for (int i = 0; i < N; ++i)
334 c[i] = (a[i] * 4) / a[i];
335
336 #pragma acc kernels async (1)
337 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
338 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
339 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
340 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
341 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
342 for (int i = 0; i < N; ++i)
343 d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
344
345 #pragma acc kernels wait (1) async (1)
346 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
347 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
348 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
349 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
350 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
351 for (int i = 0; i < N; ++i)
352 e[i] = a[i] + b[i] + c[i] + d[i];
353
354 #pragma acc wait (1)
355 }
356
357 for (int i = 0; i < N; ++i)
358 {
359 if (a[i] != 2)
360 abort ();
361 if (b[i] != 4)
362 abort ();
363 if (c[i] != 4)
364 abort ();
365 if (d[i] != 1)
366 abort ();
367 if (e[i] != 11)
368 abort ();
369 }
370
371 free (a);
372 free (b);
373 free (c);
374 free (d);
375 free (e);
376
377 return 0;
378 }