1 #include <openacc.h>
2 #include <stdlib.h>
3 #include <stdbool.h>
4
5 #define N 32
6
7 int
main(int argc,char ** argv)8 main(int argc, char **argv)
9 {
10 float *a, *b, *d_a, *d_b, exp, exp2;
11 int i;
12 const int one = 1;
13 const int zero = 0;
14 int n;
15
16 a = (float *) malloc (N * sizeof (float));
17 b = (float *) malloc (N * sizeof (float));
18 d_a = (float *) acc_malloc (N * sizeof (float));
19 d_b = (float *) acc_malloc (N * sizeof (float));
20
21 for (i = 0; i < N; i++)
22 a[i] = 4.0;
23
24 #pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(1)
25 {
26 int ii;
27
28 for (ii = 0; ii < N; ii++)
29 {
30 if (acc_on_device (acc_device_host))
31 b[ii] = a[ii] + 1;
32 else
33 b[ii] = a[ii];
34 }
35 }
36
37 #if ACC_MEM_SHARED
38 exp = 5.0;
39 #else
40 exp = 4.0;
41 #endif
42
43 for (i = 0; i < N; i++)
44 {
45 if (b[i] != exp)
46 abort();
47 }
48
49 for (i = 0; i < N; i++)
50 a[i] = 16.0;
51
52 #pragma acc parallel if(0)
53 {
54 int ii;
55
56 for (ii = 0; ii < N; ii++)
57 {
58 if (acc_on_device (acc_device_host))
59 b[ii] = a[ii] + 1;
60 else
61 b[ii] = a[ii];
62 }
63 }
64
65 for (i = 0; i < N; i++)
66 {
67 if (b[i] != 17.0)
68 abort();
69 }
70
71 for (i = 0; i < N; i++)
72 a[i] = 8.0;
73
74 #pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(one)
75 {
76 int ii;
77
78 for (ii = 0; ii < N; ii++)
79 {
80 if (acc_on_device (acc_device_host))
81 b[ii] = a[ii] + 1;
82 else
83 b[ii] = a[ii];
84 }
85 }
86
87 #if ACC_MEM_SHARED
88 exp = 9.0;
89 #else
90 exp = 8.0;
91 #endif
92
93 for (i = 0; i < N; i++)
94 {
95 if (b[i] != exp)
96 abort();
97 }
98
99 for (i = 0; i < N; i++)
100 a[i] = 22.0;
101
102 #pragma acc parallel if(zero)
103 {
104 int ii;
105
106 for (ii = 0; ii < N; ii++)
107 {
108 if (acc_on_device (acc_device_host))
109 b[ii] = a[ii] + 1;
110 else
111 b[ii] = a[ii];
112 }
113 }
114
115 for (i = 0; i < N; i++)
116 {
117 if (b[i] != 23.0)
118 abort();
119 }
120
121 for (i = 0; i < N; i++)
122 a[i] = 16.0;
123
124 #pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(true)
125 {
126 int ii;
127
128 for (ii = 0; ii < N; ii++)
129 {
130 if (acc_on_device (acc_device_host))
131 b[ii] = a[ii] + 1;
132 else
133 b[ii] = a[ii];
134 }
135 }
136
137 #if ACC_MEM_SHARED
138 exp = 17.0;
139 #else
140 exp = 16.0;
141 #endif
142
143 for (i = 0; i < N; i++)
144 {
145 if (b[i] != exp)
146 abort();
147 }
148
149 for (i = 0; i < N; i++)
150 a[i] = 76.0;
151
152 #pragma acc parallel if(false)
153 {
154 int ii;
155
156 for (ii = 0; ii < N; ii++)
157 {
158 if (acc_on_device (acc_device_host))
159 b[ii] = a[ii] + 1;
160 else
161 b[ii] = a[ii];
162 }
163 }
164
165 for (i = 0; i < N; i++)
166 {
167 if (b[i] != 77.0)
168 abort();
169 }
170
171 for (i = 0; i < N; i++)
172 a[i] = 22.0;
173
174 n = 1;
175
176 #pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(n)
177 {
178 int ii;
179
180 for (ii = 0; ii < N; ii++)
181 {
182 if (acc_on_device (acc_device_host))
183 b[ii] = a[ii] + 1;
184 else
185 b[ii] = a[ii];
186 }
187 }
188
189 #if ACC_MEM_SHARED
190 exp = 23.0;
191 #else
192 exp = 22.0;
193 #endif
194
195 for (i = 0; i < N; i++)
196 {
197 if (b[i] != exp)
198 abort();
199 }
200
201 for (i = 0; i < N; i++)
202 a[i] = 18.0;
203
204 n = 0;
205
206 #pragma acc parallel if(n)
207 {
208 int ii;
209
210 for (ii = 0; ii < N; ii++)
211 {
212 if (acc_on_device (acc_device_host))
213 b[ii] = a[ii] + 1;
214 else
215 b[ii] = a[ii];
216 }
217 }
218
219 for (i = 0; i < N; i++)
220 {
221 if (b[i] != 19.0)
222 abort();
223 }
224
225 for (i = 0; i < N; i++)
226 a[i] = 49.0;
227
228 n = 1;
229
230 #pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(n + n)
231 {
232 int ii;
233
234 for (ii = 0; ii < N; ii++)
235 {
236 if (acc_on_device (acc_device_host))
237 b[ii] = a[ii] + 1;
238 else
239 b[ii] = a[ii];
240 }
241 }
242
243 #if ACC_MEM_SHARED
244 exp = 50.0;
245 #else
246 exp = 49.0;
247 #endif
248
249 for (i = 0; i < N; i++)
250 {
251 if (b[i] != exp)
252 abort();
253 }
254
255 for (i = 0; i < N; i++)
256 a[i] = 38.0;
257
258 n = 0;
259
260 #pragma acc parallel if(n + n)
261 {
262 int ii;
263
264 for (ii = 0; ii < N; ii++)
265 {
266 if (acc_on_device (acc_device_host))
267 b[ii] = a[ii] + 1;
268 else
269 b[ii] = a[ii];
270 }
271 }
272
273 for (i = 0; i < N; i++)
274 {
275 if (b[i] != 39.0)
276 abort();
277 }
278
279 for (i = 0; i < N; i++)
280 a[i] = 91.0;
281
282 #pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(-2)
283 {
284 int ii;
285
286 for (ii = 0; ii < N; ii++)
287 {
288 if (acc_on_device (acc_device_host))
289 b[ii] = a[ii] + 1;
290 else
291 b[ii] = a[ii];
292 }
293 }
294
295 #if ACC_MEM_SHARED
296 exp = 92.0;
297 #else
298 exp = 91.0;
299 #endif
300
301 for (i = 0; i < N; i++)
302 {
303 if (b[i] != exp)
304 abort();
305 }
306
307 for (i = 0; i < N; i++)
308 a[i] = 43.0;
309
310 #pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(one == 1)
311 {
312 int ii;
313
314 for (ii = 0; ii < N; ii++)
315 {
316 if (acc_on_device (acc_device_host))
317 b[ii] = a[ii] + 1;
318 else
319 b[ii] = a[ii];
320 }
321 }
322
323 #if ACC_MEM_SHARED
324 exp = 44.0;
325 #else
326 exp = 43.0;
327 #endif
328
329 for (i = 0; i < N; i++)
330 {
331 if (b[i] != exp)
332 abort();
333 }
334
335 for (i = 0; i < N; i++)
336 a[i] = 87.0;
337
338 #pragma acc parallel if(one == 0)
339 {
340 int ii;
341
342 for (ii = 0; ii < N; ii++)
343 {
344 if (acc_on_device (acc_device_host))
345 b[ii] = a[ii] + 1;
346 else
347 b[ii] = a[ii];
348 }
349 }
350
351 for (i = 0; i < N; i++)
352 {
353 if (b[i] != 88.0)
354 abort();
355 }
356
357 for (i = 0; i < N; i++)
358 {
359 a[i] = 3.0;
360 b[i] = 9.0;
361 }
362
363 #if ACC_MEM_SHARED
364 exp = 0.0;
365 exp2 = 0.0;
366 #else
367 acc_map_data (a, d_a, N * sizeof (float));
368 acc_map_data (b, d_b, N * sizeof (float));
369 exp = 3.0;
370 exp2 = 9.0;
371 #endif
372
373 #pragma acc update device(a[0:N], b[0:N]) if(1)
374
375 for (i = 0; i < N; i++)
376 {
377 a[i] = 0.0;
378 b[i] = 0.0;
379 }
380
381 #pragma acc update host(a[0:N], b[0:N]) if(1)
382
383 for (i = 0; i < N; i++)
384 {
385 if (a[i] != exp)
386 abort();
387
388 if (b[i] != exp2)
389 abort();
390 }
391
392 for (i = 0; i < N; i++)
393 {
394 a[i] = 6.0;
395 b[i] = 12.0;
396 }
397
398 #pragma acc update device(a[0:N], b[0:N]) if(0)
399
400 for (i = 0; i < N; i++)
401 {
402 a[i] = 0.0;
403 b[i] = 0.0;
404 }
405
406 #pragma acc update host(a[0:N], b[0:N]) if(1)
407
408 for (i = 0; i < N; i++)
409 {
410 if (a[i] != exp)
411 abort();
412
413 if (b[i] != exp2)
414 abort();
415 }
416
417 for (i = 0; i < N; i++)
418 {
419 a[i] = 26.0;
420 b[i] = 21.0;
421 }
422
423 #pragma acc update device(a[0:N], b[0:N]) if(1)
424
425 for (i = 0; i < N; i++)
426 {
427 a[i] = 0.0;
428 b[i] = 0.0;
429 }
430
431 #pragma acc update host(a[0:N], b[0:N]) if(0)
432
433 for (i = 0; i < N; i++)
434 {
435 if (a[i] != 0.0)
436 abort();
437
438 if (b[i] != 0.0)
439 abort();
440 }
441
442 #if !ACC_MEM_SHARED
443 acc_unmap_data (a);
444 acc_unmap_data (b);
445 #endif
446
447 acc_free (d_a);
448 acc_free (d_b);
449
450 for (i = 0; i < N; i++)
451 {
452 a[i] = 4.0;
453 b[i] = 0.0;
454 }
455
456 #pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1)
457 {
458 #pragma acc parallel present(a[0:N])
459 {
460 int ii;
461
462 for (ii = 0; ii < N; ii++)
463 {
464 b[ii] = a[ii];
465 }
466 }
467 }
468
469 for (i = 0; i < N; i++)
470 {
471 if (b[i] != 4.0)
472 abort();
473 }
474
475 for (i = 0; i < N; i++)
476 {
477 a[i] = 8.0;
478 b[i] = 1.0;
479 }
480
481 #pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(0)
482 {
483 #if !ACC_MEM_SHARED
484 if (acc_is_present (a, N * sizeof (float)))
485 abort ();
486 #endif
487
488 #if !ACC_MEM_SHARED
489 if (acc_is_present (b, N * sizeof (float)))
490 abort ();
491 #endif
492 }
493
494 for (i = 0; i < N; i++)
495 {
496 a[i] = 18.0;
497 b[i] = 21.0;
498 }
499
500 #pragma acc data copyin(a[0:N]) if(1)
501 {
502 #if !ACC_MEM_SHARED
503 if (!acc_is_present (a, N * sizeof (float)))
504 abort ();
505 #endif
506
507 #pragma acc data copyout(b[0:N]) if(0)
508 {
509 #if !ACC_MEM_SHARED
510 if (acc_is_present (b, N * sizeof (float)))
511 abort ();
512 #endif
513
514 #pragma acc data copyout(b[0:N]) if(1)
515 {
516 #pragma acc parallel present(a[0:N]) present(b[0:N])
517 {
518 int ii;
519
520 for (ii = 0; ii < N; ii++)
521 {
522 b[ii] = a[ii];
523 }
524 }
525 }
526
527 #if !ACC_MEM_SHARED
528 if (acc_is_present (b, N * sizeof (float)))
529 abort ();
530 #endif
531 }
532 }
533
534 for (i = 0; i < N; i++)
535 {
536 if (b[i] != 18.0)
537 abort ();
538 }
539
540 #pragma acc enter data copyin (b[0:N]) if (0)
541
542 #if !ACC_MEM_SHARED
543 if (acc_is_present (b, N * sizeof (float)))
544 abort ();
545 #endif
546
547 #pragma acc exit data delete (b[0:N]) if (0)
548
549 #pragma acc enter data copyin (b[0:N]) if (1)
550
551 #if !ACC_MEM_SHARED
552 if (!acc_is_present (b, N * sizeof (float)))
553 abort ();
554 #endif
555
556 #pragma acc exit data delete (b[0:N]) if (1)
557
558 #if !ACC_MEM_SHARED
559 if (acc_is_present (b, N * sizeof (float)))
560 abort ();
561 #endif
562
563 #pragma acc enter data copyin (b[0:N]) if (zero)
564
565 #if !ACC_MEM_SHARED
566 if (acc_is_present (b, N * sizeof (float)))
567 abort ();
568 #endif
569
570 #pragma acc exit data delete (b[0:N]) if (zero)
571
572 #pragma acc enter data copyin (b[0:N]) if (one)
573
574 #if !ACC_MEM_SHARED
575 if (!acc_is_present (b, N * sizeof (float)))
576 abort ();
577 #endif
578
579 #pragma acc exit data delete (b[0:N]) if (one)
580
581 #if !ACC_MEM_SHARED
582 if (acc_is_present (b, N * sizeof (float)))
583 abort ();
584 #endif
585
586 #pragma acc enter data copyin (b[0:N]) if (one == 0)
587
588 #if !ACC_MEM_SHARED
589 if (acc_is_present (b, N * sizeof (float)))
590 abort ();
591 #endif
592
593 #pragma acc exit data delete (b[0:N]) if (one == 0)
594
595 #pragma acc enter data copyin (b[0:N]) if (one == 1)
596
597 #if !ACC_MEM_SHARED
598 if (!acc_is_present (b, N * sizeof (float)))
599 abort ();
600 #endif
601
602 #pragma acc exit data delete (b[0:N]) if (one == 1)
603
604 #if !ACC_MEM_SHARED
605 if (acc_is_present (b, N * sizeof (float)))
606 abort ();
607 #endif
608
609 for (i = 0; i < N; i++)
610 a[i] = 4.0;
611
612 #pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(1)
613 {
614 int ii;
615
616 for (ii = 0; ii < N; ii++)
617 {
618 if (acc_on_device (acc_device_host))
619 b[ii] = a[ii] + 1;
620 else
621 b[ii] = a[ii];
622 }
623 }
624
625 #if ACC_MEM_SHARED
626 exp = 5.0;
627 #else
628 exp = 4.0;
629 #endif
630
631 for (i = 0; i < N; i++)
632 {
633 if (b[i] != exp)
634 abort();
635 }
636
637 for (i = 0; i < N; i++)
638 a[i] = 16.0;
639
640 #pragma acc kernels if(0)
641 {
642 int ii;
643
644 for (ii = 0; ii < N; ii++)
645 {
646 if (acc_on_device (acc_device_host))
647 b[ii] = a[ii] + 1;
648 else
649 b[ii] = a[ii];
650 }
651 }
652
653 for (i = 0; i < N; i++)
654 {
655 if (b[i] != 17.0)
656 abort();
657 }
658
659 for (i = 0; i < N; i++)
660 a[i] = 8.0;
661
662 #pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(one)
663 {
664 int ii;
665
666 for (ii = 0; ii < N; ii++)
667 {
668 if (acc_on_device (acc_device_host))
669 b[ii] = a[ii] + 1;
670 else
671 b[ii] = a[ii];
672 }
673 }
674
675 #if ACC_MEM_SHARED
676 exp = 9.0;
677 #else
678 exp = 8.0;
679 #endif
680
681 for (i = 0; i < N; i++)
682 {
683 if (b[i] != exp)
684 abort();
685 }
686
687 for (i = 0; i < N; i++)
688 a[i] = 22.0;
689
690 #pragma acc kernels if(zero)
691 {
692 int ii;
693
694 for (ii = 0; ii < N; ii++)
695 {
696 if (acc_on_device (acc_device_host))
697 b[ii] = a[ii] + 1;
698 else
699 b[ii] = a[ii];
700 }
701 }
702
703 for (i = 0; i < N; i++)
704 {
705 if (b[i] != 23.0)
706 abort();
707 }
708
709 for (i = 0; i < N; i++)
710 a[i] = 16.0;
711
712 #pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(true)
713 {
714 int ii;
715
716 for (ii = 0; ii < N; ii++)
717 {
718 if (acc_on_device (acc_device_host))
719 b[ii] = a[ii] + 1;
720 else
721 b[ii] = a[ii];
722 }
723 }
724
725 #if ACC_MEM_SHARED
726 exp = 17.0;
727 #else
728 exp = 16.0;
729 #endif
730
731 for (i = 0; i < N; i++)
732 {
733 if (b[i] != exp)
734 abort();
735 }
736
737 for (i = 0; i < N; i++)
738 a[i] = 76.0;
739
740 #pragma acc kernels if(false)
741 {
742 int ii;
743
744 for (ii = 0; ii < N; ii++)
745 {
746 if (acc_on_device (acc_device_host))
747 b[ii] = a[ii] + 1;
748 else
749 b[ii] = a[ii];
750 }
751 }
752
753 for (i = 0; i < N; i++)
754 {
755 if (b[i] != 77.0)
756 abort();
757 }
758
759 for (i = 0; i < N; i++)
760 a[i] = 22.0;
761
762 n = 1;
763
764 #pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(n)
765 {
766 int ii;
767
768 for (ii = 0; ii < N; ii++)
769 {
770 if (acc_on_device (acc_device_host))
771 b[ii] = a[ii] + 1;
772 else
773 b[ii] = a[ii];
774 }
775 }
776
777 #if ACC_MEM_SHARED
778 exp = 23.0;
779 #else
780 exp = 22.0;
781 #endif
782
783 for (i = 0; i < N; i++)
784 {
785 if (b[i] != exp)
786 abort();
787 }
788
789 for (i = 0; i < N; i++)
790 a[i] = 18.0;
791
792 n = 0;
793
794 #pragma acc kernels if(n)
795 {
796 int ii;
797
798 for (ii = 0; ii < N; ii++)
799 {
800 if (acc_on_device (acc_device_host))
801 b[ii] = a[ii] + 1;
802 else
803 b[ii] = a[ii];
804 }
805 }
806
807 for (i = 0; i < N; i++)
808 {
809 if (b[i] != 19.0)
810 abort();
811 }
812
813 for (i = 0; i < N; i++)
814 a[i] = 49.0;
815
816 n = 1;
817
818 #pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(n + n)
819 {
820 int ii;
821
822 for (ii = 0; ii < N; ii++)
823 {
824 if (acc_on_device (acc_device_host))
825 b[ii] = a[ii] + 1;
826 else
827 b[ii] = a[ii];
828 }
829 }
830
831 #if ACC_MEM_SHARED
832 exp = 50.0;
833 #else
834 exp = 49.0;
835 #endif
836
837 for (i = 0; i < N; i++)
838 {
839 if (b[i] != exp)
840 abort();
841 }
842
843 for (i = 0; i < N; i++)
844 a[i] = 38.0;
845
846 n = 0;
847
848 #pragma acc kernels if(n + n)
849 {
850 int ii;
851
852 for (ii = 0; ii < N; ii++)
853 {
854 if (acc_on_device (acc_device_host))
855 b[ii] = a[ii] + 1;
856 else
857 b[ii] = a[ii];
858 }
859 }
860
861 for (i = 0; i < N; i++)
862 {
863 if (b[i] != 39.0)
864 abort();
865 }
866
867 for (i = 0; i < N; i++)
868 a[i] = 91.0;
869
870 #pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(-2)
871 {
872 int ii;
873
874 for (ii = 0; ii < N; ii++)
875 {
876 if (acc_on_device (acc_device_host))
877 b[ii] = a[ii] + 1;
878 else
879 b[ii] = a[ii];
880 }
881 }
882
883 #if ACC_MEM_SHARED
884 exp = 92.0;
885 #else
886 exp = 91.0;
887 #endif
888
889 for (i = 0; i < N; i++)
890 {
891 if (b[i] != exp)
892 abort();
893 }
894
895 for (i = 0; i < N; i++)
896 a[i] = 43.0;
897
898 #pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(one == 1)
899 {
900 int ii;
901
902 for (ii = 0; ii < N; ii++)
903 {
904 if (acc_on_device (acc_device_host))
905 b[ii] = a[ii] + 1;
906 else
907 b[ii] = a[ii];
908 }
909 }
910
911 #if ACC_MEM_SHARED
912 exp = 44.0;
913 #else
914 exp = 43.0;
915 #endif
916
917 for (i = 0; i < N; i++)
918 {
919 if (b[i] != exp)
920 abort();
921 }
922
923 for (i = 0; i < N; i++)
924 a[i] = 87.0;
925
926 #pragma acc kernels if(one == 0)
927 {
928 int ii;
929
930 for (ii = 0; ii < N; ii++)
931 {
932 if (acc_on_device (acc_device_host))
933 b[ii] = a[ii] + 1;
934 else
935 b[ii] = a[ii];
936 }
937 }
938
939 for (i = 0; i < N; i++)
940 {
941 if (b[i] != 88.0)
942 abort();
943 }
944
945 for (i = 0; i < N; i++)
946 {
947 a[i] = 3.0;
948 b[i] = 9.0;
949 }
950
951 #if ACC_MEM_SHARED
952 exp = 0.0;
953 exp2 = 0.0;
954 #else
955 acc_map_data (a, d_a, N * sizeof (float));
956 acc_map_data (b, d_b, N * sizeof (float));
957 exp = 3.0;
958 exp2 = 9.0;
959 #endif
960
961 return 0;
962 }
963