1 /* { dg-do run } */
2 /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
3
4 #include <openacc.h>
5 #include <string.h>
6 #include <stdio.h>
7 #include <stdlib.h>
8 #include <stdbool.h>
9
10 int
11 main (int argc, char **argv)
12 {
13 int N = 8;
14 float *a, *b, *c, *d;
15 int i;
16
17 a = (float *) malloc (N * sizeof (float));
18 b = (float *) malloc (N * sizeof (float));
19 c = (float *) malloc (N * sizeof (float));
20
21 for (i = 0; i < N; i++)
22 {
23 a[i] = 3.0;
24 b[i] = 0.0;
25 }
26
27 #pragma acc parallel copyin (a[0:N]) copyout (b[0:N])
28 {
29 int ii;
30
31 for (ii = 0; ii < N; ii++)
32 b[ii] = a[ii];
33 }
34
35 for (i = 0; i < N; i++)
36 {
37 if (b[i] != 3.0)
38 abort ();
39 }
40
41 if (acc_is_present (&a[0], (N * sizeof (float))))
42 abort ();
43
44 if (acc_is_present (&b[0], (N * sizeof (float))))
45 abort ();
46
47 for (i = 0; i < N; i++)
48 {
49 a[i] = 5.0;
50 b[i] = 1.0;
51 }
52
53 #pragma acc parallel copyin (a[0:N]) copyout (b[0:N])
54 {
55 int ii;
56
57 for (ii = 0; ii < N; ii++)
58 b[ii] = a[ii];
59 }
60
61 for (i = 0; i < N; i++)
62 {
63 if (b[i] != 5.0)
64 abort ();
65 }
66
67 if (acc_is_present (&a[0], (N * sizeof (float))))
68 abort ();
69
70 if (acc_is_present (&b[0], (N * sizeof (float))))
71 abort ();
72
73 for (i = 0; i < N; i++)
74 {
75 a[i] = 6.0;
76 b[i] = 0.0;
77 }
78
79 d = (float *) acc_copyin (&a[0], N * sizeof (float));
80
81 for (i = 0; i < N; i++)
82 {
83 a[i] = 9.0;
84 }
85
86 #pragma acc parallel present_or_copyin (a[0:N]) copyout (b[0:N])
87 {
88 int ii;
89
90 for (ii = 0; ii < N; ii++)
91 b[ii] = a[ii];
92 }
93
94 for (i = 0; i < N; i++)
95 {
96 if (b[i] != 6.0)
97 abort ();
98 }
99
100 if (!acc_is_present (&a[0], (N * sizeof (float))))
101 abort ();
102
103 if (acc_is_present (&b[0], (N * sizeof (float))))
104 abort ();
105
106 acc_delete (&a[0], N * sizeof (float));
107
108 if (acc_is_present (&a[0], N * sizeof (float)))
109 abort ();
110
111 for (i = 0; i < N; i++)
112 {
113 a[i] = 6.0;
114 b[i] = 0.0;
115 }
116
117 #pragma acc parallel copyin (a[0:N]) present_or_copyout (b[0:N])
118 {
119 int ii;
120
121 for (ii = 0; ii < N; ii++)
122 b[ii] = a[ii];
123 }
124
125 for (i = 0; i < N; i++)
126 {
127 if (b[i] != 6.0)
128 abort ();
129 }
130
131 if (acc_is_present (&a[0], (N * sizeof (float))))
132 abort ();
133
134 if (acc_is_present (&b[0], (N * sizeof (float))))
135 abort ();
136
137 for (i = 0; i < N; i++)
138 {
139 a[i] = 5.0;
140 b[i] = 2.0;
141 }
142
143 d = (float *) acc_copyin (&b[0], N * sizeof (float));
144
145 #pragma acc parallel copyin (a[0:N]) present_or_copyout (b[0:N])
146 {
147 int ii;
148
149 for (ii = 0; ii < N; ii++)
150 b[ii] = a[ii];
151 }
152
153 for (i = 0; i < N; i++)
154 {
155 if (a[i] != 5.0)
156 abort ();
157
158 if (b[i] != 2.0)
159 abort ();
160 }
161
162 if (acc_is_present (&a[0], (N * sizeof (float))))
163 abort ();
164
165 if (!acc_is_present (&b[0], (N * sizeof (float))))
166 abort ();
167
168 acc_delete (&b[0], N * sizeof (float));
169
170 if (acc_is_present (&b[0], (N * sizeof (float))))
171 abort ();
172
173 for (i = 0; i < N; i++)
174 {
175 a[i] = 3.0;
176 b[i] = 4.0;
177 }
178
179 #pragma acc parallel copy (a[0:N]) copyout (b[0:N])
180 {
181 int ii;
182
183 for (ii = 0; ii < N; ii++)
184 {
185 a[ii] = a[ii] + 1;
186 b[ii] = a[ii] + 2;
187 }
188 }
189
190 for (i = 0; i < N; i++)
191 {
192 if (a[i] != 4.0)
193 abort ();
194
195 if (b[i] != 6.0)
196 abort ();
197 }
198
199 if (acc_is_present (&a[0], (N * sizeof (float))))
200 abort ();
201
202 if (acc_is_present (&b[0], (N * sizeof (float))))
203 abort ();
204
205 for (i = 0; i < N; i++)
206 {
207 a[i] = 4.0;
208 b[i] = 7.0;
209 }
210
211 #pragma acc parallel present_or_copy (a[0:N]) present_or_copy (b[0:N])
212 {
213 int ii;
214
215 for (ii = 0; ii < N; ii++)
216 {
217 a[ii] = a[ii] + 1;
218 b[ii] = b[ii] + 2;
219 }
220 }
221
222 for (i = 0; i < N; i++)
223 {
224 if (a[i] != 5.0)
225 abort ();
226
227 if (b[i] != 9.0)
228 abort ();
229 }
230
231 if (acc_is_present (&a[0], (N * sizeof (float))))
232 abort ();
233
234 if (acc_is_present (&b[0], (N * sizeof (float))))
235 abort ();
236
237 for (i = 0; i < N; i++)
238 {
239 a[i] = 3.0;
240 b[i] = 7.0;
241 }
242
243 d = (float *) acc_copyin (&a[0], N * sizeof (float));
244 d = (float *) acc_copyin (&b[0], N * sizeof (float));
245
246 #pragma acc parallel present_or_copy (a[0:N]) present_or_copy (b[0:N])
247 {
248 int ii;
249
250 for (ii = 0; ii < N; ii++)
251 {
252 a[ii] = a[ii] + 1;
253 b[ii] = b[ii] + 2;
254 }
255 }
256
257 for (i = 0; i < N; i++)
258 {
259 if (a[i] != 3.0)
260 abort ();
261
262 if (b[i] != 7.0)
263 abort ();
264 }
265
266 if (!acc_is_present (&a[0], (N * sizeof (float))))
267 abort ();
268
269 if (!acc_is_present (&b[0], (N * sizeof (float))))
270 abort ();
271
272 acc_delete (&a[0], N * sizeof (float));
273
274 if (acc_is_present (&a[0], N * sizeof (float)))
275 abort ();
276
277 acc_delete (&b[0], N * sizeof (float));
278
279 if (acc_is_present (&b[0], N * sizeof (float)))
280 abort ();
281
282 for (i = 0; i < N; i++)
283 {
284 a[i] = 3.0;
285 b[i] = 7.0;
286 }
287
288 #pragma acc parallel copyin (a[0:N]) create (c[0:N]) copyout (b[0:N])
289 {
290 int ii;
291
292 for (ii = 0; ii < N; ii++)
293 {
294 c[ii] = a[ii];
295 b[ii] = c[ii];
296 }
297 }
298
299 for (i = 0; i < N; i++)
300 {
301 if (a[i] != 3.0)
302 abort ();
303
304 if (b[i] != 3.0)
305 abort ();
306 }
307
308 if (acc_is_present (&a[0], (N * sizeof (float))))
309 abort ();
310
311 if (acc_is_present (&b[0], (N * sizeof (float))))
312 abort ();
313
314 if (acc_is_present (&c[0], (N * sizeof (float))))
315 abort ();
316
317 for (i = 0; i < N; i++)
318 {
319 a[i] = 4.0;
320 b[i] = 8.0;
321 }
322
323 #pragma acc parallel copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N])
324 {
325 int ii;
326
327 for (ii = 0; ii < N; ii++)
328 {
329 c[ii] = a[ii];
330 b[ii] = c[ii];
331 }
332 }
333
334 for (i = 0; i < N; i++)
335 {
336 if (a[i] != 4.0)
337 abort ();
338
339 if (b[i] != 4.0)
340 abort ();
341 }
342
343 if (acc_is_present (&a[0], (N * sizeof (float))))
344 abort ();
345
346 if (acc_is_present (&b[0], (N * sizeof (float))))
347 abort ();
348
349 if (acc_is_present (&c[0], (N * sizeof (float))))
350 abort ();
351
352 for (i = 0; i < N; i++)
353 {
354 a[i] = 2.0;
355 b[i] = 5.0;
356 }
357
358 d = (float *) acc_malloc (N * sizeof (float));
359 acc_map_data (c, d, N * sizeof (float));
360
361 #pragma acc parallel copyin (a[0:N]) present_or_create (c[0:N]) copyout (b[0:N])
362 {
363 int ii;
364
365 for (ii = 0; ii < N; ii++)
366 {
367 c[ii] = a[ii];
368 b[ii] = c[ii];
369 }
370 }
371
372 for (i = 0; i < N; i++)
373 {
374 if (a[i] != 2.0)
375 abort ();
376
377 if (b[i] != 2.0)
378 abort ();
379 }
380
381 if (acc_is_present (a, (N * sizeof (float))))
382 abort ();
383
384 if (acc_is_present (b, (N * sizeof (float))))
385 abort ();
386
387 if (!acc_is_present (c, (N * sizeof (float))))
388 abort ();
389
390 d = (float *) acc_deviceptr (c);
391
392 acc_unmap_data (c);
393
394 acc_free (d);
395
396 for (i = 0; i < N; i++)
397 {
398 a[i] = 4.0;
399 b[i] = 8.0;
400 }
401
402 d = (float *) acc_malloc (N * sizeof (float));
403 acc_map_data (c, d, N * sizeof (float));
404
405 #pragma acc parallel copyin (a[0:N]) present (c[0:N]) copyout (b[0:N])
406 {
407 int ii;
408
409 for (ii = 0; ii < N; ii++)
410 {
411 c[ii] = a[ii];
412 b[ii] = c[ii];
413 }
414 }
415
416 for (i = 0; i < N; i++)
417 {
418 if (a[i] != 4.0)
419 abort ();
420
421 if (b[i] != 4.0)
422 abort ();
423 }
424
425 if (acc_is_present (a, (N * sizeof (float))))
426 abort ();
427
428 if (acc_is_present (b, (N * sizeof (float))))
429 abort ();
430
431 if (!acc_is_present (c, (N * sizeof (float))))
432 abort ();
433
434 acc_unmap_data (c);
435
436 acc_free (d);
437
438 for (i = 0; i < N; i++)
439 {
440 a[i] = 4.0;
441 b[i] = 8.0;
442 }
443
444 acc_copyin (a, N * sizeof (float));
445
446 d = (float *) acc_malloc (N * sizeof (float));
447 acc_map_data (b, d, N * sizeof (float));
448
449 d = (float *) acc_malloc (N * sizeof (float));
450 acc_map_data (c, d, N * sizeof (float));
451
452 #pragma acc parallel present (a[0:N]) present (c[0:N]) present (b[0:N])
453 {
454 int ii;
455
456 for (ii = 0; ii < N; ii++)
457 {
458 c[ii] = a[ii];
459 b[ii] = c[ii];
460 }
461 }
462
463 if (!acc_is_present (a, (N * sizeof (float))))
464 abort ();
465
466 if (!acc_is_present (b, (N * sizeof (float))))
467 abort ();
468
469 if (!acc_is_present (c, (N * sizeof (float))))
470 abort ();
471
472 d = (float *) acc_deviceptr (b);
473
474 acc_memcpy_from_device (b, d, N * sizeof (float));
475
476 for (i = 0; i < N; i++)
477 {
478 if (a[i] != 4.0)
479 abort ();
480
481 if (b[i] != 4.0)
482 abort ();
483 }
484
485 acc_delete (a, N * sizeof (float));
486
487 if (acc_is_present (a, N * sizeof (float)))
488 abort ();
489
490 d = (float *) acc_deviceptr (b);
491
492 acc_unmap_data (b);
493
494 if (acc_is_present (b, N * sizeof (float)))
495 abort ();
496
497 acc_free (d);
498
499 d = (float *) acc_deviceptr (c);
500
501 acc_unmap_data (c);
502
503 if (acc_is_present (c, N * sizeof (float)))
504 abort ();
505
506 acc_free (d);
507
508 for (i = 0; i < N; i++)
509 {
510 a[i] = 3.0;
511 b[i] = 6.0;
512 }
513
514 d = (float *) acc_malloc (N * sizeof (float));
515
516 #pragma acc parallel copyin (a[0:N]) deviceptr (d) copyout (b[0:N])
517 {
518 int ii;
519
520 for (ii = 0; ii < N; ii++)
521 {
522 d[ii] = a[ii];
523 b[ii] = d[ii];
524 }
525 }
526
527 for (i = 0; i < N; i++)
528 {
529 if (a[i] != 3.0)
530 abort ();
531
532 if (b[i] != 3.0)
533 abort ();
534 }
535
536 if (acc_is_present (a, (N * sizeof (float))))
537 abort ();
538
539 if (acc_is_present (b, (N * sizeof (float))))
540 abort ();
541
542 acc_free (d);
543
544 for (i = 0; i < N; i++)
545 {
546 a[i] = 6.0;
547 b[i] = 0.0;
548 }
549
550 d = (float *) acc_copyin (&a[0], N * sizeof (float));
551
552 for (i = 0; i < N; i++)
553 {
554 a[i] = 9.0;
555 }
556
557 #pragma acc parallel pcopyin (a[0:N]) copyout (b[0:N])
558 {
559 int ii;
560
561 for (ii = 0; ii < N; ii++)
562 b[ii] = a[ii];
563 }
564
565 for (i = 0; i < N; i++)
566 {
567 if (b[i] != 6.0)
568 abort ();
569 }
570
571 if (!acc_is_present (&a[0], (N * sizeof (float))))
572 abort ();
573
574 if (acc_is_present (&b[0], (N * sizeof (float))))
575 abort ();
576
577 acc_delete (&a[0], N * sizeof (float));
578
579 if (acc_is_present (&a[0], N * sizeof (float)))
580 abort ();
581
582 for (i = 0; i < N; i++)
583 {
584 a[i] = 6.0;
585 b[i] = 0.0;
586 }
587
588 #pragma acc parallel copyin (a[0:N]) pcopyout (b[0:N])
589 {
590 int ii;
591
592 for (ii = 0; ii < N; ii++)
593 b[ii] = a[ii];
594 }
595
596 for (i = 0; i < N; i++)
597 {
598 if (b[i] != 6.0)
599 abort ();
600 }
601
602 if (acc_is_present (&a[0], (N * sizeof (float))))
603 abort ();
604
605 if (acc_is_present (&b[0], (N * sizeof (float))))
606 abort ();
607
608 for (i = 0; i < N; i++)
609 {
610 a[i] = 6.0;
611 b[i] = 0.0;
612 }
613
614 #pragma acc parallel pcopy (a[0:N], b[0:N])
615 {
616 int ii;
617
618 for (ii = 0; ii < N; ii++)
619 b[ii] = a[ii];
620 }
621
622 for (i = 0; i < N; i++)
623 {
624 if (b[i] != 6.0)
625 abort ();
626 }
627
628 if (acc_is_present (&a[0], (N * sizeof (float))))
629 abort ();
630
631 if (acc_is_present (&b[0], (N * sizeof (float))))
632 abort ();
633
634 for (i = 0; i < N; i++)
635 {
636 a[i] = 5.0;
637 b[i] = 7.0;
638 }
639
640 #pragma acc parallel copyin (a[0:N]) pcreate (c[0:N]) copyout (b[0:N])
641 {
642 int ii;
643
644 for (ii = 0; ii < N; ii++)
645 {
646 c[ii] = a[ii];
647 b[ii] = c[ii];
648 }
649 }
650
651 for (i = 0; i < N; i++)
652 {
653 if (a[i] != 5.0)
654 abort ();
655
656 if (b[i] != 5.0)
657 abort ();
658 }
659
660 if (acc_is_present (&a[0], (N * sizeof (float))))
661 abort ();
662
663 if (acc_is_present (&b[0], (N * sizeof (float))))
664 abort ();
665
666 if (acc_is_present (&c[0], (N * sizeof (float))))
667 abort ();
668
669 return 0;
670 }