1 /* Test 'acc enter/exit data' regions with 'acc update'. */
2
3 /* { dg-do run } */
4
5 #include <stdlib.h>
6
7 int
8 main (int argc, char **argv)
9 {
10 int N = 128; //1024 * 1024;
11 float *a, *b, *c, *d, *e;
12 int i;
13 int nbytes;
14
15 nbytes = N * sizeof (float);
16
17 a = (float *) malloc (nbytes);
18 b = (float *) malloc (nbytes);
19 c = (float *) malloc (nbytes);
20 d = (float *) malloc (nbytes);
21 e = (float *) malloc (nbytes);
22
23 for (i = 0; i < N; i++)
24 {
25 a[i] = 3.0;
26 b[i] = 0.0;
27 }
28
29 #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
30 #pragma acc parallel present (a[0:N], b[0:N]) async wait
31 #pragma acc loop
32 for (i = 0; i < N; i++)
33 b[i] = a[i];
34
35 #pragma acc update host (a[0:N], b[0:N]) async wait
36 #pragma acc wait
37
38 for (i = 0; i < N; i++)
39 {
40 if (a[i] != 3.0)
41 abort ();
42
43 if (b[i] != 3.0)
44 abort ();
45 }
46
47 for (i = 0; i < N; i++)
48 {
49 a[i] = 2.0;
50 b[i] = 0.0;
51 }
52
53 #pragma acc update device (a[0:N], b[0:N]) async (1)
54 #pragma acc parallel present (a[0:N], b[0:N]) async (1)
55 #pragma acc loop
56 for (i = 0; i < N; i++)
57 b[i] = a[i];
58
59 #pragma acc update host (a[0:N], b[0:N]) async (1) wait (1)
60 #pragma acc wait (1)
61
62 for (i = 0; i < N; i++)
63 {
64 if (a[i] != 2.0)
65 abort ();
66
67 if (b[i] != 2.0)
68 abort ();
69 }
70
71 for (i = 0; i < N; i++)
72 {
73 a[i] = 3.0;
74 b[i] = 0.0;
75 c[i] = 0.0;
76 d[i] = 0.0;
77 }
78
79 #pragma acc update device (a[0:N]) async (1)
80 #pragma acc update device (b[0:N]) async (2)
81 #pragma acc enter data copyin (c[0:N], d[0:N]) async (3)
82
83 #pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1,2)
84 #pragma acc loop
85 for (i = 0; i < N; i++)
86 b[i] = (a[i] * a[i] * a[i]) / a[i];
87
88 #pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1,3)
89 #pragma acc loop
90 for (i = 0; i < N; i++)
91 c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
92
93 #pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1,3)
94 #pragma acc loop
95 for (i = 0; i < N; i++)
96 d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
97
98 #pragma acc update host (a[0:N], b[0:N], c[0:N], d[0:N]) async (1) wait (1,2,3)
99 #pragma acc wait (1)
100
101 for (i = 0; i < N; i++)
102 {
103 if (a[i] != 3.0)
104 abort ();
105
106 if (b[i] != 9.0)
107 abort ();
108
109 if (c[i] != 4.0)
110 abort ();
111
112 if (d[i] != 1.0)
113 abort ();
114 }
115
116 for (i = 0; i < N; i++)
117 {
118 a[i] = 2.0;
119 b[i] = 0.0;
120 c[i] = 0.0;
121 d[i] = 0.0;
122 e[i] = 0.0;
123 }
124
125 #pragma acc update device (a[0:N], b[0:N], c[0:N], d[0:N]) async (1)
126 #pragma acc enter data copyin (e[0:N]) async (5)
127
128 #pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1)
129 for (int ii = 0; ii < N; ii++)
130 b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
131
132 #pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1)
133 for (int ii = 0; ii < N; ii++)
134 c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
135
136 #pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1)
137 for (int ii = 0; ii < N; ii++)
138 d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
139
140 #pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \
141 wait (1, 2, 3, 5) async (4)
142 for (int ii = 0; ii < N; ii++)
143 e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
144
145 #pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
146 copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
147 #pragma acc exit data delete (N) wait(1) async(2)
148 #pragma acc wait (2)
149
150 for (i = 0; i < N; i++)
151 {
152 if (a[i] != 2.0)
153 abort ();
154
155 if (b[i] != 4.0)
156 abort ();
157
158 if (c[i] != 4.0)
159 abort ();
160
161 if (d[i] != 1.0)
162 abort ();
163
164 if (e[i] != 11.0)
165 abort ();
166 }
167
168 free (a);
169 free (b);
170 free (c);
171 free (d);
172 free (e);
173
174 return 0;
175 }