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