Bridge++  Ver.2.1.3
afield_openacc-inc.h
Go to the documentation of this file.
1 
10 //#include "inline/define_params.h"
11 //#include "inline/define_index.h"
12 
13 //====================================================================
14 void afield_init(real_t* RESTRICT data, const int size)
15 {
16 #pragma acc enter data create(data[0:size])
17 
18 #pragma acc data present(data[0:size]) copyin(size)
19 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
20  {
21 #pragma acc loop gang worker vector
22  for(int i = 0; i < size; ++i){
23  data[i] = 0.0;
24  }
25  }
26 
27 }
28 
29 //====================================================================
30 void afield_tidyup(real_t* RESTRICT data, const int size)
31 {
32 #pragma acc exit data delete(data[0:size])
33 }
34 
35 //====================================================================
36 void afield_set(real_t* RESTRICT v, real_t a, const int nin, const int nst)
37 {
38  int Nst_pad = CEIL_NWP(nst);
39  int nv = nin * Nst_pad;
40 #pragma acc data present(v[0:nv]) copyin(a, nin, nst)
41 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
42  {
43 #pragma acc loop gang worker vector
44  for(int ist = 0; ist < nst; ++ist){
45 
46  for(int in = 0; in < nin; ++in){
47  v[IDX2(nin, in, ist)] = a;
48  }
49 
50  }
51  }
52 
53 }
54 
55 //====================================================================
56 void copy_to_device(real_t* RESTRICT v, int nv)
57 {
58 #pragma acc update device(v[0:nv])
59 }
60 
61 //====================================================================
62 void copy_to_device(real_t* RESTRICT v, int nv1, int nv)
63 {
64 #pragma acc update device(v[nv1:nv])
65 }
66 
67 //====================================================================
68 void copy_from_device(real_t* RESTRICT v, int nv)
69 {
70 #pragma acc update host(v[0:nv])
71 }
72 
73 //====================================================================
74 void copy_from_device(real_t* RESTRICT v, int nv1, int nv)
75 {
76 
77 #pragma acc update host(v[nv1:nv])
78 }
79 
80 //====================================================================
81 void convert(real_t* RESTRICT v, double* RESTRICT w,
82  int nin, int nvol, int nvol_pad)
83 {
84  // w is assumed to be Field, thus its volume is not ceiled.
85  int nv1 = nin * nvol;
86  int nv2 = nin * nvol_pad;
87 
88 #pragma acc data present(v[0:nv2]) copyin(w[0:nv1], nin, nvol, nvol_pad)
89 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
90  {
91 #pragma acc loop gang worker vector
92  for(int ist = 0; ist < nvol_pad; ++ist){
93  if(ist < nvol){
94  for(int in = 0; in < nin; ++in){
95  v[IDX2(nin, in, ist)] = w[IDX_CORE(nin, in, ist)];
96  }
97  }else{
98  for(int in = 0; in < nin; ++in){
99  v[IDX2(nin, in, ist)] = 0.0;
100  }
101  }
102  }
103  }
104 
105 }
106 
107 //====================================================================
108 void reverse(double* RESTRICT v, real_t* RESTRICT w,
109  int nin, int nvol, int nvol_pad)
110 {
111  // v is assumed to be Field, thus its volume is not ceiled.
112  int nv1 = nin * nvol;
113  int nv2 = nin * nvol_pad;
114 
115 #pragma acc data present(w[0:nv2]) copyin(nin, nvol, nvol_pad) \
116  copyout(v[0:nv1])
117 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
118  {
119 #pragma acc loop gang worker vector
120  for(int ist = 0; ist < nvol; ++ist){
121  for(int in = 0; in < nin; ++in){
122  v[IDX_CORE(nin, in, ist)] = w[IDX2(nin, in, ist)];
123  }
124  }
125  }
126 
127 }
128 
129 //====================================================================
130 void copy(real_t* RESTRICT v, real_t* RESTRICT w, int nin, int nvol)
131 {
132  int Nvol_pad = CEIL_NWP(nvol);
133  int nv = nin * Nvol_pad;
134 
135 #pragma acc data present(v[0:nv], w[0:nv]) copyin(nin, nvol)
136 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
137  {
138 #pragma acc loop gang worker vector
139  for(int ist = 0; ist < nvol; ++ist){
140  for(int in = 0; in < nin; ++in){
141  v[IDX2(nin, in, ist)] = w[IDX2(nin, in, ist)];
142  }
143  }
144  }
145 
146 }
147 
148 //====================================================================
149 void copy(real_t* RESTRICT v, int nv1, real_t* RESTRICT w, int nv2,
150  int nin, int nvol)
151 {
152  int Nvol_pad = CEIL_NWP(nvol);
153  int nv = nin * Nvol_pad;
154 
155 #pragma acc data present(v[nv1:nv], w[nv2:nv]) \
156  copyin(nin, nvol, nv1, nv2)
157 
158 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
159  {
160 #pragma acc loop gang worker vector
161  for(int ist = 0; ist < nvol; ++ist){
162  for(int in = 0; in < nin; ++in){
163  v[nv1 + IDX2(nin, in, ist)] = w[nv2 + IDX2(nin, in, ist)];
164  }
165  }
166  }
167 
168 }
169 
170 //====================================================================
171 void axpy(real_t *RESTRICT v, int nv1, real_t a,
172  real_t *RESTRICT w, int nv2, int nin, int nvol)
173 {
174  int Nvol_pad = CEIL_NWP(nvol);
175  int nv = nin * Nvol_pad;
176 
177 #pragma acc data present(v[nv1:nv],w[nv2:nv]) \
178  copyin(a, nin, nvol, nv1, nv2)
179 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
180  {
181 #pragma acc loop gang worker vector
182  for(int ist = 0; ist < nvol; ++ist){
183  for(int in = 0; in < nin; ++in){
184  v[nv1 + IDX2(nin, in, ist)] += a * w[nv2 + IDX2(nin, in, ist)];
185  }
186  }
187  }
188 
189 }
190 
191 //====================================================================
192 void axpy(real_t *RESTRICT v, int nv1, real_t ar, real_t ai,
193  real_t *RESTRICT w, int nv2, int nin, int nvol)
194 {
195  int Nvol_pad = CEIL_NWP(nvol);
196  int nv = nin * Nvol_pad;
197 
198 #pragma acc data present(v[nv1:nv],w[nv2:nv]) \
199  copyin(ar, ai, nin, nvol, nv1, nv2)
200 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
201  {
202  int nin2 = nin/2;
203 
204 #pragma acc loop gang worker vector
205  for(int ist = 0; ist < nvol; ++ist){
206  for(int in = 0; in < nin2; ++in){
207  real_t wr = w[nv2 + IDX2(nin, 2*in, ist)];
208  real_t wi = w[nv2 + IDX2(nin, 2*in+1, ist)];
209  v[nv1 + IDX2(nin, 2*in, ist)] += ar * wr - ai * wi;
210  v[nv1 + IDX2(nin, 2*in+1, ist)] += ai * wr + ar * wi;
211  }
212  }
213  }
214 
215 }
216 
217 //====================================================================
218 void aypx(real_t a, real_t *RESTRICT v, int nv1,
219  real_t *RESTRICT w, int nv2, int nin, int nvol)
220 {
221  int Nvol_pad = CEIL_NWP(nvol);
222  int nv = nin * Nvol_pad;
223 
224 #pragma acc data present(v[nv1:nv],w[nv2:nv])
225 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
226  {
227 #pragma acc loop gang worker vector
228  for(int ist = 0; ist < nvol; ++ist){
229  for(int in = 0; in < nin; ++in){
230  v[nv1 + IDX2(nin, in, ist)] = a * v[nv1 + IDX2(nin, in, ist)]
231  + w[nv2 + IDX2(nin, in, ist)];
232  }
233  }
234  }
235 
236 }
237 
238 //====================================================================
239 void aypx(real_t ar, real_t ai, real_t *RESTRICT v, int nv1,
240  real_t *RESTRICT w, int nv2, int nin, int nvol)
241 {
242  int Nvol_pad = CEIL_NWP(nvol);
243  int nv = nin * Nvol_pad;
244 
245 #pragma acc data present(v[nv1:nv],w[nv2:nv]) \
246  copyin(ar, ai, nin, nvol)
247 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
248  {
249  int nin2 = nin/2;
250 #pragma acc loop gang worker vector
251  for(int ist = 0; ist < nvol; ++ist){
252  for(int in = 0; in < nin2; ++in){
253  real_t vr = v[nv1 + IDX2(nin, 2*in, ist)];
254  real_t vi = v[nv1 + IDX2(nin, 2*in+1, ist)];
255  v[nv1 + IDX2(nin, 2*in, ist)] = ar * vr - ai * vi
256  + w[nv2 + IDX2(nin, 2*in, ist)];
257  v[nv1 + IDX2(nin, 2*in+1, ist)] = ai * vr + ar * vi
258  + w[nv2 + IDX2(nin, 2*in+1, ist)];
259  }
260  }
261 
262  }
263 
264 }
265 
266 //====================================================================
267 void scal(real_t* RESTRICT v, int nv1, real_t a, int nin, int nvol)
268 {
269  int Nvol_pad = CEIL_NWP(nvol);
270  int nv = nin * Nvol_pad;
271 
272 #pragma acc data present(v[nv1:nv]) copyin(a, nin, nvol)
273 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
274  {
275 #pragma acc loop gang worker vector
276  for(int ist = 0; ist < nvol; ++ist){
277  for(int in = 0; in < nin; ++in){
278  v[nv1 + IDX2(nin, in, ist)] *= a;
279  }
280  }
281  }
282 
283 }
284 
285 //====================================================================
286 void scal(real_t* RESTRICT v, int nv1, real_t ar, real_t ai,
287  int nin, int nvol)
288 {
289  int Nvol_pad = CEIL_NWP(nvol);
290  int nv = nin * Nvol_pad;
291  int nin2 = nin/2;
292 
293 #pragma acc data present(v[nv1:nv]) copyin(ar, ai, nin, nvol)
294 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
295  {
296 #pragma acc loop gang worker vector
297  for(int ist = 0; ist < nvol; ++ist){
298  for(int in = 0; in < nin2; ++in){
299  real_t vr = v[nv1 + IDX2(nin, 2*in, ist)];
300  real_t vi = v[nv1 + IDX2(nin, 2*in+1, ist)];
301  v[nv1 + IDX2(nin, 2*in, ist)] = ar * vr - ai * vi;
302  v[nv1 + IDX2(nin, 2*in+1, ist)] = ai * vr + ar * vi;
303  }
304  }
305 
306  }
307 
308 }
309 
310 //====================================================================
311 real_t dot(real_t* RESTRICT v1, real_t* RESTRICT v2,
312  int nin, int nvol)
313 {
314  int Nvol_pad = CEIL_NWP(nvol);
315  int nv = nin * Nvol_pad;
316 
317  real_t a = 0.0;
318 
319 #pragma acc data present(v1[0:nv], v2[0:nv]) copyin(nin, nvol)
320 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
321  {
322 #pragma acc loop gang worker vector reduction(+:a)
323  for(int ist = 0; ist < nvol; ++ist){
324  for(int in = 0; in < nin; ++in){
325  a += v1[IDX2(nin, in, ist)] * v2[IDX2(nin, in, ist)];
326  }
327  }
328  }
329 
330  return a;
331 
332 }
333 
334 //====================================================================
335 real_t norm2(real_t* RESTRICT v1, int nin, int nvol){
336 
337  int Nvol_pad = CEIL_NWP(nvol);
338  int nv = nin * Nvol_pad;
339 
340  double a = 0.0;
341 
342 #pragma acc data present(v1[0:nv])
343 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
344  {
345 #pragma acc loop gang worker vector reduction(+:a)
346  for(int ist = 0; ist < nvol; ++ist){
347  for(int in = 0; in < nin; ++in){
348  a += double(v1[IDX2(nin, in, ist)] * v1[IDX2(nin, in, ist)]);
349  }
350  }
351  }
352 
353  return real_t(a);
354 
355 }
356 
357 //====================================================================
358 void dotc(real_t* ar, real_t* ai,
359  real_t *RESTRICT v1, real_t *RESTRICT v2,
360  int nin, int nst)
361 {
362 
363  double ar2 = 0.0;
364  double ai2 = 0.0;
365 
366  int Nst_pad = CEIL_NWP(nst);
367  int nv = nin * Nst_pad;
368 
369 #pragma acc data present(v1[0:nv], v2[0:nv]) copyin(nst, nin)
370 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
371  {
372  int nin2 = nin/2;
373 
374 #pragma acc loop gang worker vector reduction(+:ar2,ai2)
375  for(int ist = 0; ist < nst; ++ist){
376 
377  double ar1 = 0.0;
378  double ai1 = 0.0;
379  for(int in = 0; in < nin2; ++in){
380  real_t v1r = v1[IDX2(nin, 2*in, ist)];
381  real_t v1i = v1[IDX2(nin, 2*in+1, ist)];
382  real_t v2r = v2[IDX2(nin, 2*in, ist)];
383  real_t v2i = v2[IDX2(nin, 2*in+1, ist)];
384  ar1 += double(v1r * v2r + v1i * v2i);
385  ai1 += double(v1r * v2i - v1i * v2r);
386  }
387  ar2 += ar1;
388  ai2 += ai1;
389  }
390 
391  } // acc parallel
392 
393  *ar = real_t(ar2);
394  *ai = real_t(ai2);
395 
396 }
397 
398 //====================================================================
399 void xI(real_t *RESTRICT v1, int nin, int nst){
400 
401  int Nst_pad = CEIL_NWP(nst);
402  int nv = nin * Nst_pad;
403 
404 #pragma acc data present(v1[0:nv]) copyin(nin, nst)
405 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
406  {
407  int nin2 = nin/2;
408 
409 #pragma acc loop gang worker vector
410  for(int ist = 0; ist < nst; ++ist){
411  for(int in = 0; in < nin2; ++in){
412  real_t vr = v1[IDX2(nin, 2*in, ist)];
413  real_t vi = v1[IDX2(nin, 2*in+1, ist)];
414  v1[IDX2(nin, 2*in, ist)] = -vi;
415  v1[IDX2(nin, 2*in+1, ist)] = vr;
416  }
417  }
418 
419  }
420 
421 }
422 
423 //====================================================================
424 void conjg(real_t *RESTRICT v1, int nin, int nst){
425 
426  int Nst_pad = CEIL_NWP(nst);
427  int nv = nin * Nst_pad;
428 
429 #pragma acc data present(v1[0:nv]) copyin(nin, nst)
430 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
431  {
432  int nin2 = nin/2;
433 
434 #pragma acc loop gang worker vector
435  for(int ist = 0; ist < nst; ++ist){
436  for(int in = 0; in < nin2; ++in){
437  real_t vi = v1[IDX2(nin, 2*in+1, ist)];
438  v1[IDX2(nin, 2*in+1, ist)] = -vi;
439  }
440  }
441 
442  }
443 
444 }
445 
446 //====================================================================
scal
void scal(real_t *RESTRICT v, int nv1, real_t a, int nin, int nvol)
Definition: afield_openacc-inc.h:267
copy_to_device
void copy_to_device(real_t *RESTRICT v, int nv)
Definition: afield_openacc-inc.h:56
afield_set
void afield_set(real_t *RESTRICT v, real_t a, const int nin, const int nst)
Definition: afield_openacc-inc.h:36
IDX_CORE
#define IDX_CORE(nin, in, ist)
Definition: define_index.h:25
convert
void convert(real_t *RESTRICT v, double *RESTRICT w, int nin, int nvol, int nvol_pad)
Definition: afield_openacc-inc.h:81
conjg
void conjg(real_t *RESTRICT v1, int nin, int nst)
Definition: afield_openacc-inc.h:424
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
reverse
void reverse(double *RESTRICT v, real_t *RESTRICT w, int nin, int nvol, int nvol_pad)
Definition: afield_openacc-inc.h:108
copy_from_device
void copy_from_device(real_t *RESTRICT v, int nv)
Definition: afield_openacc-inc.h:68
dot
real_t dot(real_t *RESTRICT v1, real_t *RESTRICT v2, int nin, int nvol)
Definition: afield_openacc-inc.h:311
afield_tidyup
void afield_tidyup(real_t *RESTRICT data, const int size)
Definition: afield_openacc-inc.h:30
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
afield_init
void afield_init(real_t *RESTRICT data, const int size)
Definition: afield_openacc-inc.h:14
axpy
void axpy(real_t *RESTRICT v, int nv1, real_t a, real_t *RESTRICT w, int nv2, int nin, int nvol)
Definition: afield_openacc-inc.h:171
IDX2
#define IDX2(nin, in, ist)
Definition: define_index.h:28
aypx
void aypx(real_t a, real_t *RESTRICT v, int nv1, real_t *RESTRICT w, int nv2, int nin, int nvol)
Definition: afield_openacc-inc.h:218
norm2
real_t norm2(real_t *RESTRICT v1, int nin, int nvol)
Definition: afield_openacc-inc.h:335
xI
void xI(real_t *RESTRICT v1, int nin, int nst)
Definition: afield_openacc-inc.h:399
copy
void copy(real_t *RESTRICT v, real_t *RESTRICT w, int nin, int nvol)
Definition: afield_openacc-inc.h:130
dotc
void dotc(real_t *ar, real_t *ai, real_t *RESTRICT v1, real_t *RESTRICT v2, int nin, int nst)
Definition: afield_openacc-inc.h:358