Bridge++  Ver.2.1.3
shiftAField_lex_openacc-inc.h
Go to the documentation of this file.
1 
9 #include "lib_alt_Accel/inline/define_params.h"
10 #include "lib_alt_Accel/inline/define_index.h"
11 
12 //====================================================================
13 void shift_lex_xp1(real_t *RESTRICT buf, real_t *RESTRICT v1,
14  int nin, int *Nsize, int *bc)
15 {
16  // int idir = 0;
17 
18  int Nx = Nsize[0];
19  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
20 
21  real_t bc2 = bc[0];
22 
23  int size = nin * CEIL_NWP(Nx * Nyzt);
24  int size_b = nin * CEIL_NWP(Nyzt);
25 
26 #pragma acc data present(v1[0:size], buf[0:size_b]), copyin(bc2, Nx, Nyzt)
27 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
28  {
29 
30 #pragma acc loop gang worker vector
31  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
32  int ix = 0;
33  int ist = ix + Nx * iyzt;
34  for(int in = 0; in < nin; ++in){
35  buf[in + nin * iyzt] = bc2 * v1[IDX2(nin, in, ist)];
36  }
37  }
38 
39  }
40 
41 }
42 
43 //====================================================================
44 void shift_lex_xp2(real_t *RESTRICT v2, real_t *RESTRICT buf,
45  int nin, int *Nsize, int *bc)
46 {
47  // int idir = 0;
48 
49  int Nx = Nsize[0];
50  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
51 
52  int size = nin * CEIL_NWP(Nx * Nyzt);
53  int size_b = nin * CEIL_NWP(Nyzt);
54 
55 #pragma acc data present(v2[0:size], buf[0:size_b]), copyin(Nx, Nyzt)
56 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
57  {
58 
59 #pragma acc loop gang worker vector
60  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
61  int ix = Nx-1;
62  int ist = ix + Nx * iyzt;
63  for(int in = 0; in < nin; ++in){
64  v2[IDX2(nin, in, ist)] = buf[in + nin * iyzt];
65  }
66  }
67 
68  }
69 
70 }
71 
72 //====================================================================
73 void shift_lex_xpb(real_t *RESTRICT v2, real_t *RESTRICT v1,
74  int nin, int *Nsize, int *bc)
75 {
76  // int idir = 0;
77 
78  int Nx = Nsize[0];
79  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
80  int Nst = Nx * Nyzt;
81 
82  int size = nin * CEIL_NWP(Nx * Nyzt);
83 
84 #pragma acc data present(v2[0:size], v1[0:size]), \
85  copyin(bc[0:4], Nx, Nyzt, Nst)
86 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
87  {
88 
89 #pragma acc loop gang worker vector
90  for(int ist = 0; ist < Nst; ++ist){
91  int ix = ist % Nx;
92  int iyzt = ist/Nx;
93  int nei = ((ix+1) % Nx) + Nx * iyzt;
94  real_t bc2 = 1.0;
95  if(ix == Nx-1) bc2 = bc[0];
96 
97  for(int in = 0; in < nin; ++in){
98  v2[IDX2(nin, in, ist)] = bc2 * v1[IDX2(nin, in, nei)];
99  }
100 
101  }
102 
103  }
104 
105 }
106 
107 //====================================================================
108 void shift_lex_xm1(real_t *RESTRICT buf, real_t *RESTRICT v1,
109  int nin, int *Nsize, int *bc)
110 {
111  // int idir = 0;
112 
113  int Nx = Nsize[0];
114  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
115 
116  int size = nin * CEIL_NWP(Nx * Nyzt);
117  int size_b = nin * CEIL_NWP(Nyzt);
118 
119 #pragma acc data present(v1[0:size], buf[0:size_b]), copyin(Nx, Nyzt)
120 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
121  {
122 
123 #pragma acc loop gang worker vector
124  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
125  int ix = Nx-1;
126  int ist = ix + Nx * iyzt;
127  for(int in = 0; in < nin; ++in){
128  buf[in + nin * iyzt] = v1[IDX2(nin, in, ist)];
129  }
130  }
131 
132  }
133 
134 }
135 
136 //====================================================================
137 void shift_lex_xm2(real_t *RESTRICT v2, real_t *RESTRICT buf,
138  int nin, int *Nsize, int *bc)
139 {
140  // int idir = 0;
141 
142  int Nx = Nsize[0];
143  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
144 
145  real_t bc2 = bc[0];
146 
147  int size = nin * CEIL_NWP(Nx * Nyzt);
148  int size_b = nin * CEIL_NWP(Nyzt);
149 
150 #pragma acc data present(v2[0:size], buf[0:size_b]), copyin(bc2, Nx, Nyzt)
151 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
152  {
153 
154 #pragma acc loop gang worker vector
155  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
156  int ix = 0;
157  int ist = ix + Nx * iyzt;
158  for(int in = 0; in < nin; ++in){
159  v2[IDX2(nin, in, ist)] = bc2 * buf[in + nin * iyzt];
160  }
161  }
162 
163  }
164 
165 }
166 
167 //====================================================================
168 void shift_lex_xmb(real_t *RESTRICT v2, real_t *RESTRICT v1,
169  int nin, int *Nsize, int *bc)
170 {
171  // int idir = 0;
172 
173  int Nx = Nsize[0];
174  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
175  int Nst = Nx * Nyzt;
176 
177  int size = nin * CEIL_NWP(Nx * Nyzt);
178 
179 #pragma acc data present(v2[0:size], v1[0:size]), \
180  copyin(bc[0:4], Nx, Nyzt, Nst)
181 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
182  {
183 
184 #pragma acc loop gang worker vector
185  for(int ist = 0; ist < Nst; ++ist){
186  int ix = ist % Nx;
187  int iyzt = ist/Nx;
188  int nei = ix-1 + Nx * iyzt;
189  if(ix == 0) nei = Nx-1 + Nx * iyzt;
190  real_t bc2 = 1.0;
191  if(ix == 0) bc2 = bc[0];
192 
193  for(int in = 0; in < nin; ++in){
194  v2[IDX2(nin, in, ist)] = bc2 * v1[IDX2(nin, in, nei)];
195  }
196 
197  }
198 
199  }
200 
201 }
202 
203 //====================================================================
204 void shift_lex_yp1(real_t *RESTRICT buf, real_t *RESTRICT v1,
205  int nin, int *Nsize, int *bc)
206 {
207  int Nx = Nsize[0];
208  int Ny = Nsize[1];
209  int Nzt = Nsize[2] * Nsize[3];
210 
211  real_t bc2 = bc[1];
212 
213  int size = nin * CEIL_NWP(Nx * Ny * Nzt);
214  int size_b = nin * CEIL_NWP(Nx * Nzt);
215 
216 #pragma acc data present(v1[0:size], buf[0:size_b]), \
217  copyin(bc2, Nx, Ny, Nzt)
218 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
219  {
220 
221 #pragma acc loop gang worker vector collapse(2)
222  for(int izt = 0; izt < Nzt; ++izt){
223  for(int ix = 0; ix < Nx; ++ix){
224  int iy = 0;
225  int ist = ix + Nx * (iy + Ny * izt);
226  int ixzt = ix + Nx * izt;
227  for(int in = 0; in < nin; ++in){
228  buf[in + nin * ixzt] = bc2 * v1[IDX2(nin, in, ist)];
229  }
230  }
231  }
232 
233  }
234 
235 }
236 
237 //====================================================================
238 void shift_lex_yp2(real_t *RESTRICT v2, real_t *RESTRICT buf,
239  int nin, int *Nsize, int *bc)
240 {
241  int idir = 1;
242 
243  int Nx = Nsize[0];
244  int Ny = Nsize[1];
245  int Nzt = Nsize[2] * Nsize[3];
246  int Nst = Nx * Ny * Nzt;
247 
248  int size = nin * CEIL_NWP(Nx * Ny * Nzt);
249  int size_b = nin * CEIL_NWP(Nx * Nzt);
250 
251 #pragma acc data present(v2[0:size], buf[0:size_b]), \
252  copyin(Nx, Ny, Nzt, Nst, idir)
253 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
254  {
255 
256 #pragma acc loop gang worker vector collapse(2)
257  for(int izt = 0; izt < Nzt; ++izt){
258  for(int ix = 0; ix < Nx; ++ix){
259  int iy = Ny-1;
260  int ist = ix + Nx * (iy + Ny * izt);
261  int ixzt = ix + Nx * izt;
262 
263  for(int in = 0; in < nin; ++in){
264  v2[IDX2(nin, in, ist)] = buf[in + nin * ixzt];
265  }
266  }
267  }
268 
269  }
270 
271 }
272 
273 //====================================================================
274 void shift_lex_ypb(real_t *RESTRICT v2, real_t *RESTRICT v1,
275  int nin, int *Nsize, int *bc)
276 {
277  int idir = 1;
278 
279  int Nx = Nsize[0];
280  int Ny = Nsize[1];
281  int Nzt = Nsize[2] * Nsize[3];
282  int Nst = Nx * Ny * Nzt;
283 
284  int size = nin * CEIL_NWP(Nx * Ny * Nzt);
285 
286 #pragma acc data present(v2[0:size], v1[0:size]), \
287  copyin(bc[0:4], idir, Nx, Ny, Nzt, Nst)
288 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
289  {
290 
291 #pragma acc loop gang worker vector
292  for(int ist = 0; ist < Nst; ++ist){
293  int ix = ist % Nx;
294  int iy = (ist/Nx) % Ny;
295  int izt = ist/(Nx*Ny);
296  int iy2 = (iy+1) % Ny;
297  int nei = ix + Nx * (iy2 + Ny * izt);
298  real_t bc2 = 1.0;
299  if(iy == Ny-1) bc2 = bc[idir];
300 
301  for(int in = 0; in < nin; ++in){
302  v2[IDX2(nin, in, ist)] = bc2 * v1[IDX2(nin, in, nei)];
303  }
304  }
305 
306  }
307 
308 }
309 
310 //====================================================================
311 void shift_lex_ym1(real_t *RESTRICT buf, real_t *RESTRICT v1,
312  int nin, int *Nsize, int *bc)
313 {
314  // int idir = 1;
315 
316  int Nx = Nsize[0];
317  int Ny = Nsize[1];
318  int Nzt = Nsize[2] * Nsize[3];
319 
320  int size = nin * CEIL_NWP(Nx * Ny * Nzt);
321  int size_b = nin * CEIL_NWP(Nx * Nzt);
322 
323 #pragma acc data present(v1[0:size], buf[0:size_b]), \
324  copyin(Nx, Ny, Nzt)
325 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
326  {
327 
328 #pragma acc loop gang worker vector collapse(2)
329  for(int izt = 0; izt < Nzt; ++izt){
330  for(int ix = 0; ix < Nx; ++ix){
331  int iy = Ny-1;
332  int ist = ix + Nx * (iy + Ny*izt);
333  int ixzt = ix + Nx * izt;
334 
335  for(int in = 0; in < nin; ++in){
336  buf[in + nin * ixzt] = v1[IDX2(nin, in, ist)];
337  }
338 
339  }
340  }
341 
342  }
343 
344 }
345 
346 //====================================================================
347 void shift_lex_ym2(real_t *RESTRICT v2, real_t *RESTRICT buf,
348  int nin, int *Nsize, int *bc)
349 {
350  // int idir = 1;
351 
352  int Nx = Nsize[0];
353  int Ny = Nsize[1];
354  int Nzt = Nsize[2] * Nsize[3];
355 
356  real_t bc2 = bc[1];
357 
358  int size = nin * CEIL_NWP(Nx * Ny * Nzt);
359  int size_b = nin * CEIL_NWP(Nx * Nzt);
360 
361 #pragma acc data present(v2[0:size], buf[0:size_b]), \
362  copyin(bc2, Nx, Ny, Nzt)
363 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
364  {
365 
366 #pragma acc loop gang worker vector collapse(2)
367  for(int izt = 0; izt < Nzt; ++izt){
368  for(int ix = 0; ix < Nx; ++ix){
369  int iy = 0;
370  int ist = ix + Nx*(iy + Ny*izt);
371  int ixzt = ix + Nx * izt;
372 
373  for(int in = 0; in < nin; ++in){
374  v2[IDX2(nin, in, ist)] = bc2 * buf[in + nin*ixzt];
375  }
376 
377  }
378  }
379 
380  }
381 
382 }
383 
384 //====================================================================
385 void shift_lex_ymb(real_t *RESTRICT v2, real_t *RESTRICT v1,
386  int nin, int *Nsize, int *bc)
387 {
388  int idir = 1;
389 
390  int Nx = Nsize[0];
391  int Ny = Nsize[1];
392  int Nzt = Nsize[2] * Nsize[3];
393  int Nst = Nx * Ny * Nzt;
394 
395  int size = nin * CEIL_NWP(Nx * Ny * Nzt);
396 
397 #pragma acc data present(v2[0:size], v1[0:size]), \
398  copyin(bc[0:4], idir, Nx, Ny, Nzt, Nst)
399 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
400  {
401 
402 #pragma acc loop gang worker vector
403  for(int ist = 0; ist < Nst; ++ist){
404  int ix = ist % Nx;
405  int iy = (ist/Nx) % Ny;
406  int izt = ist/(Nx*Ny);
407  int iy2 = (iy-1+Ny) % Ny;
408  int nei = ix + Nx * (iy2 + Ny * izt);
409  real_t bc2 = 1.0;
410  if(iy == 0) bc2 = bc[idir];
411 
412  for(int in = 0; in < nin; ++in){
413  v2[IDX2(nin, in, ist)] = bc2 * v1[IDX2(nin, in, nei)];
414  }
415 
416  }
417 
418  }
419 
420 }
421 
422 //====================================================================
423 void shift_lex_zp1(real_t *RESTRICT buf, real_t *RESTRICT v1,
424  int nin, int *Nsize, int *bc)
425 {
426  int idir = 2;
427 
428  int Nxy = Nsize[0] * Nsize[1];
429  int Nz = Nsize[2];
430  int Nt = Nsize[3];
431 
432  real_t bc2 = bc[idir];
433 
434  int size = nin * CEIL_NWP(Nxy * Nz * Nt);
435  int size_b = nin * CEIL_NWP(Nxy * Nt);
436 
437 #pragma acc data present(v1[0:size], buf[0:size_b]), \
438  copyin(bc2, Nxy, Nz, Nt)
439 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
440  {
441 
442 #pragma acc loop gang worker vector collapse(2)
443  for(int it = 0; it < Nt; ++it){
444  for(int ixy = 0; ixy < Nxy; ++ixy){
445  int iz = 0;
446  int ist = ixy + Nxy * (iz + Nz * it);
447  int ixyt = ixy + Nxy * it;
448 
449  for(int in = 0; in < nin; ++in){
450  buf[in + nin * ixyt] = bc2 * v1[IDX2(nin, in, ist)];
451  }
452 
453  }
454  }
455 
456  }
457 
458 }
459 
460 //====================================================================
461 void shift_lex_zp2(real_t *RESTRICT v2, real_t *RESTRICT buf,
462  int nin, int *Nsize, int *bc)
463 {
464  int idir = 2;
465 
466  int Nxy = Nsize[0] * Nsize[1];
467  int Nz = Nsize[2];
468  int Nt = Nsize[3];
469  int Nst = Nxy * Nz * Nt;
470 
471  int size = nin * CEIL_NWP(Nxy * Nz * Nt);
472  int size_b = nin * CEIL_NWP(Nxy * Nt);
473 
474 #pragma acc data present(v2[0:size], buf[0:size_b]), \
475  copyin(Nxy, Nz, Nt, Nst, idir)
476 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
477  {
478 
479 #pragma acc loop gang worker vector collapse(2)
480  for(int it = 0; it < Nt; ++it){
481  for(int ixy = 0; ixy < Nxy; ++ixy){
482  int iz = Nz-1;
483  int ist = ixy + Nxy * (iz + Nz * it);
484  int ixyt = ixy + Nxy * it;
485 
486  for(int in = 0; in < nin; ++in){
487  v2[IDX2(nin, in, ist)] = buf[in + nin * ixyt];
488  }
489 
490  }
491  }
492 
493  }
494 
495 }
496 
497 //====================================================================
498 void shift_lex_zpb(real_t *RESTRICT v2, real_t *RESTRICT v1,
499  int nin, int *Nsize, int *bc)
500 {
501  int idir = 2;
502 
503  int Nxy = Nsize[0] * Nsize[1];
504  int Nz = Nsize[2];
505  int Nt = Nsize[3];
506  int Nst = Nxy * Nz * Nt;
507 
508  int size = nin * CEIL_NWP(Nxy * Nz * Nt);
509 
510 #pragma acc data present(v2[0:size], v1[0:size]), \
511  copyin(bc[0:4], idir, Nxy, Nz, Nt, Nst)
512 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
513  {
514 
515 #pragma acc loop gang worker vector
516  for(int ist = 0; ist < Nst; ++ist){
517  int ixy = ist % Nxy;
518  int iz = (ist/Nxy) % Nz;
519  int it = ist/(Nxy*Nz);
520  int iz2 = (iz+1) % Nz;
521  int nei = ixy + Nxy * (iz2 + Nz * it);
522  real_t bc2 = 1.0;
523  if(iz == Nz-1) bc2 = bc[idir];
524 
525  for(int in = 0; in < nin; ++in){
526  v2[IDX2(nin, in, ist)] = bc2 * v1[IDX2(nin, in, nei)];
527  }
528 
529  }
530 
531  }
532 
533 }
534 
535 //====================================================================
536 void shift_lex_zm1(real_t *RESTRICT buf, real_t *RESTRICT v1,
537  int nin, int *Nsize, int *bc)
538 {
539  int idir = 2;
540 
541  int Nxy = Nsize[0] * Nsize[1];
542  int Nz = Nsize[2];
543  int Nt = Nsize[3];
544  int Nst = Nxy * Nz * Nt;
545 
546  int size = nin * CEIL_NWP(Nxy * Nz * Nt);
547  int size_b = nin * CEIL_NWP(Nxy * Nt);
548 
549 #pragma acc data present(v1[0:size], buf[0:size_b]), \
550  copyin(Nxy, Nz, Nt, Nst, idir)
551 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
552  {
553 
554 #pragma acc loop gang worker vector collapse(2)
555  for(int it = 0; it < Nt; ++it){
556  for(int ixy = 0; ixy < Nxy; ++ixy){
557  int iz = Nz-1;
558  int ist = ixy + Nxy * (iz + Nz * it);
559  int ixyt = ixy + Nxy * it;
560 
561  for(int in = 0; in < nin; ++in){
562  buf[in + nin * ixyt] = v1[IDX2(nin, in, ist)];
563  }
564  }
565  }
566 
567  }
568 
569 }
570 
571 //====================================================================
572 void shift_lex_zm2(real_t *RESTRICT v2, real_t *RESTRICT buf,
573  int nin, int *Nsize, int *bc)
574 {
575  int idir = 2;
576 
577  int Nxy = Nsize[0] * Nsize[1];
578  int Nz = Nsize[2];
579  int Nt = Nsize[3];
580 
581  real_t bc2 = bc[idir];
582 
583  int size = nin * CEIL_NWP(Nxy * Nz * Nt);
584  int size_b = nin * CEIL_NWP(Nxy * Nt);
585 
586 #pragma acc data present(v2[0:size], buf[0:size_b]), \
587  copyin(bc2, Nxy, Nz, Nt)
588 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
589  {
590 
591 #pragma acc loop gang worker vector collapse(2)
592  for(int it = 0; it < Nt; ++it){
593  for(int ixy = 0; ixy < Nxy; ++ixy){
594  int iz = 0;
595  int ist = ixy + Nxy * (iz + Nz * it);
596  int ixyt = ixy + Nxy * it;
597 
598  for(int in = 0; in < nin; ++in){
599  v2[IDX2(nin, in, ist)] = bc2 * buf[in + nin*ixyt];
600  }
601 
602  }
603  }
604 
605  }
606 
607 }
608 
609 //====================================================================
610 void shift_lex_zmb(real_t *RESTRICT v2, real_t *RESTRICT v1,
611  int nin, int *Nsize, int *bc)
612 {
613  int idir = 2;
614 
615  int Nxy = Nsize[0] * Nsize[1];
616  int Nz = Nsize[2];
617  int Nt = Nsize[3];
618  int Nst = Nxy * Nz * Nt;
619 
620  int size = nin * CEIL_NWP(Nxy * Nz * Nt);
621 
622 #pragma acc data present(v2[0:size], v1[0:size]), \
623  copyin(bc[0:4], idir, Nxy, Nz, Nt, Nst)
624 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
625  {
626 
627 #pragma acc loop gang worker vector
628  for(int ist = 0; ist < Nst; ++ist){
629  int ixy = ist % Nxy;
630  int iz = (ist/Nxy) % Nz;
631  int it = ist/(Nxy*Nz);
632  int iz2 = (iz-1+Nz) % Nz;
633  int nei = ixy + Nxy * (iz2 + Nz * it);
634  real_t bc2 = 1.0;
635  if(iz == 0) bc2 = bc[idir];
636 
637  for(int in = 0; in < nin; ++in){
638  v2[IDX2(nin, in, ist)] = bc2 * v1[IDX2(nin, in, nei)];
639  }
640 
641  }
642 
643  }
644 
645 }
646 
647 //====================================================================
648 void shift_lex_tp1(real_t *RESTRICT buf, real_t *RESTRICT v1,
649  int nin, int *Nsize, int *bc)
650 {
651  int idir = 3;
652 
653  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
654  int Nt = Nsize[3];
655 
656  real_t bc2 = bc[idir];
657 
658  int size = nin * CEIL_NWP(Nxyz * Nt);
659  int size_b = nin * CEIL_NWP(Nxyz);
660 
661 #pragma acc data present(v1[0:size], buf[0:size_b]), copyin(bc2, Nxyz)
662 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
663  {
664 
665 #pragma acc loop gang worker vector
666  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
667  int it = 0;
668  int ist = ixyz + Nxyz * it;
669  for(int in = 0; in < nin; ++in){
670  buf[in + nin * ixyz] = bc2 * v1[IDX2(nin, in, ist)];
671  }
672  }
673 
674  }
675 
676 }
677 
678 //====================================================================
679 void shift_lex_tp2(real_t *RESTRICT v2, real_t *RESTRICT buf,
680  int nin, int *Nsize, int *bc)
681 {
682  int idir = 3;
683 
684  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
685  int Nt = Nsize[3];
686  int Nst = Nxyz * Nt;
687 
688  int size = nin * CEIL_NWP(Nxyz * Nt);
689  int size_b = nin * CEIL_NWP(Nxyz);
690 
691 #pragma acc data present(v2[0:size], buf[0:size_b]), \
692  copyin(Nxyz, Nst, idir)
693 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
694  {
695 
696 #pragma acc loop gang worker vector
697  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
698  int it = Nt-1;
699  int ist = ixyz + Nxyz * it;
700  for(int in = 0; in < nin; ++in){
701  v2[IDX2(nin, in, ist)] = buf[in + nin * ixyz];
702  }
703  }
704 
705  }
706 
707 }
708 
709 //====================================================================
710 void shift_lex_tpb(real_t *RESTRICT v2, real_t *RESTRICT v1,
711  int nin, int *Nsize, int *bc)
712 {
713  int idir = 3;
714 
715  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
716  int Nt = Nsize[3];
717  int Nst = Nxyz * Nt;
718 
719  int size = nin * CEIL_NWP(Nxyz * Nt);
720 
721 #pragma acc data present(v2[0:size], v1[0:size]), \
722  copyin(bc[0:4], idir, Nxyz, Nt, Nst)
723 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
724  {
725 
726 #pragma acc loop gang worker vector
727  for(int ist = 0; ist < Nst; ++ist){
728  int ixyz = ist % Nxyz;
729  int it = ist/Nxyz;
730  int it2 = (it+1) % Nt;
731  int nei = ixyz + Nxyz * it2;
732  real_t bc2 = 1.0;
733  if(it == Nt-1) bc2 = bc[idir];
734 
735  for(int in = 0; in < nin; ++in){
736  v2[IDX2(nin, in, ist)] = bc2 * v1[IDX2(nin, in, nei)];
737  }
738 
739  }
740 
741  }
742 
743 }
744 
745 //====================================================================
746 void shift_lex_tm1(real_t *RESTRICT buf, real_t *RESTRICT v1,
747  int nin, int *Nsize, int *bc)
748 {
749  int idir = 3;
750 
751  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
752  int Nt = Nsize[3];
753  int Nst = Nxyz * Nt;
754 
755  int size = nin * CEIL_NWP(Nxyz * Nt);
756  int size_b = nin * CEIL_NWP(Nxyz);
757 
758 #pragma acc data present(v1[0:size], buf[0:size_b]), \
759  copyin(Nxyz, Nt, Nst, idir)
760 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
761  {
762 
763 #pragma acc loop gang worker vector
764  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
765  int it = Nt-1;
766  int ist = ixyz + Nxyz * it;
767  for(int in = 0; in < nin; ++in){
768  buf[in + nin * ixyz] = v1[IDX2(nin, in, ist)];
769  }
770  }
771 
772  }
773 
774 }
775 
776 //====================================================================
777 void shift_lex_tm2(real_t *RESTRICT v2, real_t *RESTRICT buf,
778  int nin, int *Nsize, int *bc)
779 {
780  int idir = 3;
781 
782  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
783  int Nt = Nsize[3];
784 
785  real_t bc2 = bc[idir];
786 
787  int size = nin * CEIL_NWP(Nxyz * Nt);
788  int size_b = nin * CEIL_NWP(Nxyz);
789 
790 #pragma acc data present(v2[0:size], buf[0:size_b]), copyin(bc2, Nxyz)
791 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
792  {
793 
794 #pragma acc loop gang worker vector
795  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
796  int it = 0;
797  int ist = ixyz + Nxyz * it;
798  for(int in = 0; in < nin; ++in){
799  v2[IDX2(nin, in, ist)] = bc2 * buf[in + nin * ixyz];
800  }
801  }
802 
803  }
804 
805 }
806 
807 //====================================================================
808 void shift_lex_tmb(real_t *RESTRICT v2, real_t *RESTRICT v1,
809  int nin, int *Nsize, int *bc)
810 {
811  int idir = 3;
812 
813  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
814  int Nt = Nsize[3];
815  int Nst = Nxyz * Nt;
816 
817  int size = nin * CEIL_NWP(Nxyz * Nt);
818 
819 #pragma acc data present(v2[0:size], v1[0:size]), \
820  copyin(bc[0:4], idir, Nxyz, Nt, Nst)
821 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
822  {
823 
824 #pragma acc loop gang worker vector
825  for(int ist = 0; ist < Nst; ++ist){
826  int ixyz = ist % Nxyz;
827  int it = ist/Nxyz;
828  int it2 = (it-1+Nt) % Nt;
829  int nei = ixyz + Nxyz * it2;
830  real_t bc2 = 1.0;
831  if(it == 0) bc2 = bc[idir];
832 
833  for(int in = 0; in < nin; ++in){
834  v2[IDX2(nin, in, ist)] = bc2 * v1[IDX2(nin, in, nei)];
835  }
836  }
837 
838  }
839 
840 }
841 
842 //============================================================END=====
843 
shift_lex_tm1
void shift_lex_tm1(real_t *RESTRICT buf, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:746
shift_lex_tp1
void shift_lex_tp1(real_t *RESTRICT buf, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:648
shift_lex_ym2
void shift_lex_ym2(real_t *RESTRICT v2, real_t *RESTRICT buf, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:347
iy
int iy
Definition: mult_Wilson_xyz_openacc-inc.h:239
ixy
int ixy
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:513
izt
int izt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:262
shift_lex_zp1
void shift_lex_zp1(real_t *RESTRICT buf, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:423
shift_lex_zmb
void shift_lex_zmb(real_t *RESTRICT v2, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:610
shift_lex_tpb
void shift_lex_tpb(real_t *RESTRICT v2, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:710
shift_lex_zm1
void shift_lex_zm1(real_t *RESTRICT buf, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:536
shift_lex_xmb
void shift_lex_xmb(real_t *RESTRICT v2, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:168
shift_lex_tp2
void shift_lex_tp2(real_t *RESTRICT v2, real_t *RESTRICT buf, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:679
shift_lex_ypb
void shift_lex_ypb(real_t *RESTRICT v2, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:274
shift_lex_xm2
void shift_lex_xm2(real_t *RESTRICT v2, real_t *RESTRICT buf, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:137
shift_lex_tmb
void shift_lex_tmb(real_t *RESTRICT v2, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:808
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
iyzt
int iyzt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:14
ix
int ix
Definition: mult_Wilson_xyz_openacc-inc.h:16
shift_lex_xpb
void shift_lex_xpb(real_t *RESTRICT v2, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:73
idir
idir
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:264
shift_lex_xp1
void shift_lex_xp1(real_t *RESTRICT buf, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:13
bc2
int bc2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:62
shift_lex_yp1
void shift_lex_yp1(real_t *RESTRICT buf, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:204
shift_lex_ym1
void shift_lex_ym1(real_t *RESTRICT buf, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:311
shift_lex_zm2
void shift_lex_zm2(real_t *RESTRICT v2, real_t *RESTRICT buf, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:572
it
int it
Definition: mult_Wilson_xyz_openacc-inc.h:461
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
IDX2
#define IDX2(nin, in, ist)
Definition: define_index.h:28
define_index.h
shift_lex_zpb
void shift_lex_zpb(real_t *RESTRICT v2, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:498
shift_lex_xp2
void shift_lex_xp2(real_t *RESTRICT v2, real_t *RESTRICT buf, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:44
shift_lex_ymb
void shift_lex_ymb(real_t *RESTRICT v2, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:385
iz
int iz
Definition: mult_Wilson_xyz_openacc-inc.h:462
shift_lex_zp2
void shift_lex_zp2(real_t *RESTRICT v2, real_t *RESTRICT buf, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:461
ixyz
int ixyz
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:13
define_params.h
shift_lex_xm1
void shift_lex_xm1(real_t *RESTRICT buf, real_t *RESTRICT v1, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:108
shift_lex_tm2
void shift_lex_tm2(real_t *RESTRICT v2, real_t *RESTRICT buf, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:777
shift_lex_yp2
void shift_lex_yp2(real_t *RESTRICT v2, real_t *RESTRICT buf, int nin, int *Nsize, int *bc)
Definition: shiftAField_lex_openacc-inc.h:238