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 int NDIV2 = N / 2;
15 float *a, *b, *c;
16 float *d_a, *d_b, *d_c;
17 int i;
18
19 a = (float *) malloc (N * sizeof (float));
20 b = (float *) malloc (N * sizeof (float));
21 c = (float *) malloc (N * sizeof (float));
22
23 d_a = (float *) acc_malloc (N * sizeof (float));
24 d_b = (float *) acc_malloc (N * sizeof (float));
25 d_c = (float *) acc_malloc (N * sizeof (float));
26
27 for (i = 0; i < N; i++)
28 {
29 a[i] = 3.0;
30 b[i] = 0.0;
31 }
32
33 acc_map_data (a, d_a, N * sizeof (float));
34 acc_map_data (b, d_b, N * sizeof (float));
35 acc_map_data (c, d_c, N * sizeof (float));
36
37 #pragma acc update device (a[0:N], b[0:N])
38
39 #pragma acc parallel present (a[0:N], b[0:N])
40 {
41 int ii;
42
43 for (ii = 0; ii < N; ii++)
44 b[ii] = a[ii];
45 }
46
47 #pragma acc update host (a[0:N], b[0:N])
48
49 for (i = 0; i < N; i++)
50 {
51 if (a[i] != 3.0)
52 abort ();
53
54 if (b[i] != 3.0)
55 abort ();
56 }
57
58 if (!acc_is_present (&a[0], (N * sizeof (float))))
59 abort ();
60
61 if (!acc_is_present (&b[0], (N * sizeof (float))))
62 abort ();
63
64 for (i = 0; i < N; i++)
65 {
66 a[i] = 5.0;
67 b[i] = 1.0;
68 }
69
70 #pragma acc update device (a[0:N], b[0:N])
71
72 #pragma acc parallel present (a[0:N], b[0:N])
73 {
74 int ii;
75
76 for (ii = 0; ii < N; ii++)
77 b[ii] = a[ii];
78 }
79
80 #pragma acc update host (a[0:N], b[0:N])
81
82 for (i = 0; i < N; i++)
83 {
84 if (a[i] != 5.0)
85 abort ();
86
87 if (b[i] != 5.0)
88 abort ();
89 }
90
91 if (!acc_is_present (&a[0], (N * sizeof (float))))
92 abort ();
93
94 if (!acc_is_present (&b[0], (N * sizeof (float))))
95 abort ();
96
97 for (i = 0; i < N; i++)
98 {
99 a[i] = 5.0;
100 b[i] = 1.0;
101 }
102
103 #pragma acc update device (a[0:N], b[0:N])
104
105 #pragma acc parallel present (a[0:N], b[0:N])
106 {
107 int ii;
108
109 for (ii = 0; ii < N; ii++)
110 b[ii] = a[ii];
111 }
112
113 #pragma acc update host (a[0:N], b[0:N])
114
115 for (i = 0; i < N; i++)
116 {
117 if (a[i] != 5.0)
118 abort ();
119
120 if (b[i] != 5.0)
121 abort ();
122 }
123
124 if (!acc_is_present (&a[0], (N * sizeof (float))))
125 abort ();
126
127 if (!acc_is_present (&b[0], (N * sizeof (float))))
128 abort ();
129
130 for (i = 0; i < N; i++)
131 {
132 a[i] = 6.0;
133 b[i] = 0.0;
134 }
135
136 #pragma acc update device (a[0:N], b[0:N])
137
138 for (i = 0; i < N; i++)
139 {
140 a[i] = 9.0;
141 }
142
143 #pragma acc parallel present (a[0:N], b[0:N])
144 {
145 int ii;
146
147 for (ii = 0; ii < N; ii++)
148 b[ii] = a[ii];
149 }
150
151 #pragma acc update host (a[0:N], b[0:N])
152
153 for (i = 0; i < N; i++)
154 {
155 if (a[i] != 6.0)
156 abort ();
157
158 if (b[i] != 6.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 for (i = 0; i < N; i++)
169 {
170 a[i] = 7.0;
171 b[i] = 2.0;
172 }
173
174 #pragma acc update device (a[0:N], b[0:N])
175
176 for (i = 0; i < N; i++)
177 {
178 a[i] = 9.0;
179 }
180
181 #pragma acc parallel present (a[0:N], b[0:N])
182 {
183 int ii;
184
185 for (ii = 0; ii < N; ii++)
186 b[ii] = a[ii];
187 }
188
189 #pragma acc update host (a[0:N], b[0:N])
190
191 for (i = 0; i < N; i++)
192 {
193 if (a[i] != 7.0)
194 abort ();
195
196 if (b[i] != 7.0)
197 abort ();
198 }
199
200 for (i = 0; i < N; i++)
201 {
202 a[i] = 9.0;
203 }
204
205 #pragma acc update device (a[0:N])
206
207 #pragma acc parallel present (a[0:N], b[0:N])
208 {
209 int ii;
210
211 for (ii = 0; ii < N; ii++)
212 b[ii] = a[ii];
213 }
214
215 #pragma acc update host (a[0:N], b[0:N])
216
217 for (i = 0; i < N; i++)
218 {
219 if (a[i] != 9.0)
220 abort ();
221
222 if (b[i] != 9.0)
223 abort ();
224 }
225
226 if (!acc_is_present (&a[0], (N * sizeof (float))))
227 abort ();
228
229 if (!acc_is_present (&b[0], (N * sizeof (float))))
230 abort ();
231
232 for (i = 0; i < N; i++)
233 {
234 a[i] = 5.0;
235 }
236
237 #pragma acc update device (a[0:N])
238
239 for (i = 0; i < N; i++)
240 {
241 a[i] = 6.0;
242 }
243
244 #pragma acc update device (a[0:NDIV2])
245
246 #pragma acc parallel present (a[0:N], b[0:N])
247 {
248 int ii;
249
250 for (ii = 0; ii < N; ii++)
251 b[ii] = a[ii];
252 }
253
254 #pragma acc update host (a[0:N], b[0:N])
255
256 for (i = 0; i < NDIV2; i++)
257 {
258 if (a[i] != 6.0)
259 abort ();
260
261 if (b[i] != 6.0)
262 abort ();
263 }
264
265 for (i = NDIV2; i < N; i++)
266 {
267 if (a[i] != 5.0)
268 abort ();
269
270 if (b[i] != 5.0)
271 abort ();
272 }
273
274 if (!acc_is_present (&a[0], (N * sizeof (float))))
275 abort ();
276
277 if (!acc_is_present (&b[0], (N * sizeof (float))))
278 abort ();
279
280 for (i = 0; i < N; i++)
281 {
282 a[i] = 0.0;
283 }
284
285 #pragma acc update device (a[0:4])
286
287 #pragma acc parallel present (a[0:N])
288 {
289 int ii;
290
291 for (ii = 0; ii < N; ii++)
292 a[ii] = a[ii] + 1.0;
293 }
294
295 #pragma acc update host (a[4:4])
296
297 for (i = 0; i < NDIV2; i++)
298 {
299 if (a[i] != 0.0)
300 abort ();
301 }
302
303 for (i = NDIV2; i < N; i++)
304 {
305 if (a[i] != 6.0)
306 abort ();
307 }
308
309 #pragma acc update host (a[0:4])
310
311 for (i = 0; i < NDIV2; i++)
312 {
313 if (a[i] != 1.0)
314 abort ();
315 }
316
317 for (i = NDIV2; i < N; i++)
318 {
319 if (a[i] != 6.0)
320 abort ();
321 }
322
323 a[2] = 9;
324 a[3] = 9;
325 a[4] = 9;
326 a[5] = 9;
327
328 #pragma acc update device (a[2:4])
329
330 #pragma acc parallel present (a[0:N])
331 {
332 int ii;
333
334 for (ii = 0; ii < N; ii++)
335 a[ii] = a[ii] + 1.0;
336 }
337
338 #pragma acc update host (a[2:4])
339
340 for (i = 0; i < 2; i++)
341 {
342 if (a[i] != 1.0)
343 abort ();
344 }
345
346 for (i = 2; i < 6; i++)
347 {
348 if (a[i] != 10.0)
349 abort ();
350 }
351
352 for (i = 6; i < N; i++)
353 {
354 if (a[i] != 6.0)
355 abort ();
356 }
357
358 return 0;
359 }