1 /* acc_present_or_create, acc_present_or_copyin, etc. */
2 /* See also Fortran variants in "../libgomp.oacc-fortran/lib-32*". */
3
4 #include <stdbool.h>
5 #include <stdlib.h>
6 #include <openacc.h>
7
8 int
9 main (int argc, char **argv)
10 {
11 int *h, *d;
12 const int N = 10000;
13 const int S = N * sizeof *h;
14 bool shared_mem;
15
16 h = (int *) malloc (S);
17 if (!h)
18 abort ();
19 for (int i = 0; i < N; ++i)
20 h[i] = i + 0;
21
22 shared_mem = acc_is_present (h, S);
23
24 d = (int *) acc_present_or_create (h, S);
25 if (!d)
26 abort ();
27 if (shared_mem)
28 if (h != d)
29 abort ();
30 if (!acc_is_present (h, S))
31 abort ();
32
33 #pragma acc parallel loop deviceptr (d)
34 for (int i = 0; i < N; ++i)
35 {
36 d[i] = i + 1;
37 }
38
39 for (int i = 0; i < N; ++i)
40 {
41 if (h[i] != i + (shared_mem ? 1 : 0))
42 abort ();
43 h[i] = i + 2;
44 }
45
46 {
47 int *d_ = (int *) acc_present_or_create (h, S);
48 if (d_ != d)
49 abort ();
50 }
51
52 #pragma acc parallel loop deviceptr (d)
53 for (int i = 0; i < N; ++i)
54 {
55 if (d[i] != i + (shared_mem ? 2 : 1))
56 abort ();
57 d[i] = i + 3;
58 }
59
60 for (int i = 0; i < N; ++i)
61 {
62 if (h[i] != i + (shared_mem ? 3 : 2))
63 abort ();
64 h[i] = i + 4;
65 }
66
67 {
68 int *d_ = (int *) acc_pcreate (h, S);
69 if (d_ != d)
70 abort ();
71 }
72
73 #pragma acc parallel loop deviceptr (d)
74 for (int i = 0; i < N; ++i)
75 {
76 if (d[i] != i + (shared_mem ? 4 : 3))
77 abort ();
78 d[i] = i + 5;
79 }
80
81 for (int i = 0; i < N; ++i)
82 {
83 if (h[i] != i + (shared_mem ? 5 : 4))
84 abort ();
85 h[i] = i + 6;
86 }
87
88 {
89 int *d_ = (int *) acc_present_or_copyin (h, S);
90 if (d_ != d)
91 abort ();
92 }
93
94 #pragma acc parallel loop deviceptr (d)
95 for (int i = 0; i < N; ++i)
96 {
97 if (d[i] != i + (shared_mem ? 6 : 5))
98 abort ();
99 d[i] = i + 7;
100 }
101
102 for (int i = 0; i < N; ++i)
103 {
104 if (h[i] != i + (shared_mem ? 7 : 6))
105 abort ();
106 h[i] = i + 8;
107 }
108
109 {
110 int *d_ = (int *) acc_pcopyin (h, S);
111 if (d_ != d)
112 abort ();
113 }
114
115 #pragma acc parallel loop deviceptr (d)
116 for (int i = 0; i < N; ++i)
117 {
118 if (d[i] != i + (shared_mem ? 8 : 7))
119 abort ();
120 d[i] = i + 9;
121 }
122
123 for (int i = 0; i < N; ++i)
124 {
125 if (h[i] != i + (shared_mem ? 9 : 8))
126 abort ();
127 h[i] = i + 10;
128 }
129
130 acc_copyout_finalize (h, S);
131 d = NULL;
132 if (!shared_mem)
133 if (acc_is_present (h, S))
134 abort ();
135
136 for (int i = 0; i < N; ++i)
137 {
138 if (h[i] != i + (shared_mem ? 10 : 9))
139 abort ();
140 }
141
142 d = (int *) acc_pcopyin (h, S);
143 if (!d)
144 abort ();
145 if (shared_mem)
146 if (h != d)
147 abort ();
148 if (!acc_is_present (h, S))
149 abort ();
150
151 #pragma acc parallel loop deviceptr (d)
152 for (int i = 0; i < N; ++i)
153 {
154 if (d[i] != i + (shared_mem ? 10 : 9))
155 abort ();
156 d[i] = i + 11;
157 }
158
159 for (int i = 0; i < N; ++i)
160 {
161 if (h[i] != i + (shared_mem ? 11 : 9))
162 abort ();
163 h[i] = i + 12;
164 }
165
166 {
167 int *d_ = (int *) acc_pcopyin (h, S);
168 if (d_ != d)
169 abort ();
170 }
171
172 #pragma acc parallel loop deviceptr (d)
173 for (int i = 0; i < N; ++i)
174 {
175 if (d[i] != i + (shared_mem ? 12 : 11))
176 abort ();
177 d[i] = i + 13;
178 }
179
180 for (int i = 0; i < N; ++i)
181 {
182 if (h[i] != i + (shared_mem ? 13 : 12))
183 abort ();
184 h[i] = i + 14;
185 }
186
187 {
188 int *d_ = (int *) acc_pcreate (h, S);
189 if (d_ != d)
190 abort ();
191 }
192
193 #pragma acc parallel loop deviceptr (d)
194 for (int i = 0; i < N; ++i)
195 {
196 if (d[i] != i + (shared_mem ? 14 : 13))
197 abort ();
198 d[i] = i + 15;
199 }
200
201 for (int i = 0; i < N; ++i)
202 {
203 if (h[i] != i + (shared_mem ? 15 : 14))
204 abort ();
205 h[i] = i + 16;
206 }
207
208 {
209 int *d_ = (int *) acc_pcreate (h, S);
210 if (d_ != d)
211 abort ();
212 }
213
214 #pragma acc parallel loop deviceptr (d)
215 for (int i = 0; i < N; ++i)
216 {
217 if (d[i] != i + (shared_mem ? 16 : 15))
218 abort ();
219 d[i] = i + 17;
220 }
221
222 for (int i = 0; i < N; ++i)
223 {
224 if (h[i] != i + (shared_mem ? 17 : 16))
225 abort ();
226 h[i] = i + 18;
227 }
228
229 acc_update_self (h, S);
230 if (!acc_is_present (h, S))
231 abort ();
232
233 for (int i = 0; i < N; ++i)
234 {
235 if (h[i] != i + (shared_mem ? 18 : 17))
236 abort ();
237 }
238
239 acc_delete_finalize (h, S);
240 d = NULL;
241 if (!shared_mem)
242 if (acc_is_present (h, S))
243 abort();
244
245 free (h);
246
247 return 0;
248 }