Bridge++  Ver.2.1.3
mult_Domainwall_5din_dir_openacc-inc.h
Go to the documentation of this file.
1 
10 #ifndef MULT_DOMAINWALL_5DIN_DIR_ACC_INCLUDED
11 #define MULT_DOMAINWALL_5DIN_DIR_ACC_INCLUDED
12 
13 //====================================================================
14 void mult_domainwall_5din_xpb(
15  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
16  int Ns, int *bc, int *Nsize, int *do_comm, int flag)
17 {
18  int Nx = Nsize[0];
19  int Ny = Nsize[1];
20  int Nz = Nsize[2];
21  int Nt = Nsize[3];
22  int Nst = Nx * Ny * Nz * Nt;
23  int Nst_pad = CEIL_NWP(Nst);
24 
25  int Nin5 = NVCD * Ns;
26 
27  int size = Nin5 * Nst_pad;
28  int size_u = NDF * Nst_pad * NDIM;
29 
30 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
31  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
32  bc[0:NDIM], do_comm[0:NDIM])
33  {
34 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
35  {
36 
37 #pragma acc loop gang worker vector
38  for (int site = 0; site < Nst; ++site) {
39  int ix = site % Nx;
40  int iyzt = site / Nx;
41 
42  int idir;
43 
44  for(int is = 0; is < Ns; ++is){
45 
46  real_t vL[NVCD];
47 
48  if(flag == 0){
49  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
50  vL[ivcd] = 0.0;
51  }
52  }else{
53  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
54  vL[ivcd] = vp[IDX2(Nin5, (ivcd + NVCD * is), site)];
55  }
56  }
57 
58  idir = 0;
59 
60  if ((ix < Nx-1) || (do_comm[idir] == 0)) {
61  int ix2 = (ix + 1) % Nx;
62  int nei = ix2 + Nx * iyzt;
63  real_t bc2 = 1.0;
64  if(ix == Nx-1) bc2 = bc[0];
65 
66  real_t ut[NDF];
67  load_u(ut, up, site + Nst_pad * idir);
68 
69  real_t wt[NVCD];
70  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
71  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
72  }
73  mult_wilson_xpb(vL, ut, wt);
74  }
75 
76  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
77  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
78  }
79 
80  } // is loop
81 
82  } // site loop
83  } // acc parallel
84  } // acc data
85 
86 }
87 
88 
89 //====================================================================
90 void mult_domainwall_5din_xmb(
91  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
92  int Ns, int *bc, int *Nsize, int *do_comm, int flag)
93 {
94  int Nx = Nsize[0];
95  int Ny = Nsize[1];
96  int Nz = Nsize[2];
97  int Nt = Nsize[3];
98  int Nst = Nx * Ny * Nz * Nt;
99  int Nst_pad = CEIL_NWP(Nst);
100 
101  int Nin5 = NVCD * Ns;
102 
103  int size = Nin5 * Nst_pad;
104  int size_u = NDF * Nst_pad * NDIM;
105 
106 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
107  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns,\
108  bc[0:NDIM], do_comm[0:NDIM])
109  {
110 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
111  {
112 
113 #pragma acc loop gang worker vector
114  for (int site = 0; site < Nst; ++site) {
115 
116  int ix = site % Nx;
117  int iyzt = site / Nx;
118 
119  int idir;
120 
121  for(int is = 0; is < Ns; ++is){
122 
123  real_t vL[NVCD];
124 
125  if(flag == 0){
126  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
127  vL[ivcd] = 0.0;
128  }
129  }else{
130  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
131  vL[ivcd] = vp[IDX2(Nin5, (ivcd + NVCD * is), site)];
132  }
133  }
134 
135  idir = 0;
136 
137  if ((ix > 0) || (do_comm[idir] == 0)) {
138  int ix2 = (ix - 1 + Nx) % Nx;
139  int nei = ix2 + Nx * iyzt;
140  real_t bc2 = 1.0;
141  if(ix == 0) bc2 = bc[0];
142 
143  real_t ut[NDF];
144  load_u(ut, up, nei + Nst_pad * idir);
145 
146  real_t wt[NVCD];
147  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
148  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
149  }
150  mult_wilson_xmb(vL, ut, wt);
151  }
152 
153  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
154  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
155  }
156 
157  } // is loop
158 
159  } // site loop
160  } // acc parallel
161  } // acc data
162 
163 }
164 
165 
166 //====================================================================
167 void mult_domainwall_5din_ypb(
168  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
169  int Ns, int *bc, int *Nsize, int *do_comm, int flag)
170 {
171  int Nx = Nsize[0];
172  int Ny = Nsize[1];
173  int Nz = Nsize[2];
174  int Nt = Nsize[3];
175  int Nst = Nx * Ny * Nz * Nt;
176  int Nst_pad = CEIL_NWP(Nst);
177 
178  int Nin5 = NVCD * Ns;
179 
180  int size = Nin5 * Nst_pad;
181  int size_u = NDF * Nst_pad * NDIM;
182 
183 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
184  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
185  bc[0:NDIM], do_comm[0:NDIM])
186  {
187 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
188  {
189 
190 #pragma acc loop gang worker vector
191  for (int site = 0; site < Nst; ++site) {
192 
193  int ix = site % Nx;
194  int iyzt = site / Nx;
195  int iy = iyzt % Ny;
196  int izt = iyzt / Ny;
197 
198  for(int is = 0; is < Ns; ++is){
199 
200  real_t vL[NVCD];
201 
202  if(flag == 0){
203  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
204  vL[ivcd] = 0.0;
205  }
206  }else{
207  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
208  vL[ivcd] = vp[IDX2(Nin5, (ivcd + NVCD * is), site)];
209  }
210  }
211 
212  int idir = 1;
213 
214  if ((iy < Ny-1) || (do_comm[idir] == 0)) {
215  int iy2 = (iy + 1) % Ny;
216  int nei = ix + Nx * (iy2 + Ny * izt);
217  real_t bc2 = 1.0;
218  if(iy == Ny-1) bc2 = bc[idir];
219 
220  real_t ut[NDF];
221  load_u(ut, up, site + Nst_pad * idir);
222 
223  real_t wt[NVCD];
224  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
225  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
226  }
227  mult_wilson_ypb(vL, ut, wt);
228  }
229 
230 
231  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
232  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
233  }
234 
235  } // is loop
236 
237  } // site loop
238  } // acc parallel
239  } // acc data
240 
241 }
242 
243 
244 //====================================================================
245 void mult_domainwall_5din_ymb(
246  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
247  int Ns, int *bc, int *Nsize, int *do_comm, int flag)
248 {
249  int Nx = Nsize[0];
250  int Ny = Nsize[1];
251  int Nz = Nsize[2];
252  int Nt = Nsize[3];
253  int Nst = Nx * Ny * Nz * Nt;
254  int Nst_pad = CEIL_NWP(Nst);
255 
256  int Nin5 = NVCD * Ns;
257 
258  int size = Nin5 * Nst_pad;
259  int size_u = NDF * Nst_pad * NDIM;
260 
261 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
262  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
263  bc[0:NDIM], do_comm[0:NDIM])
264  {
265 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
266  {
267 
268 #pragma acc loop gang worker vector
269  for (int site = 0; site < Nst; ++site) {
270  int Nxy = Nx * Ny;
271 
272  int ix = site % Nx;
273  int iyzt = site / Nx;
274  int iy = iyzt % Ny;
275  int izt = site / Nxy;
276 
277  int idir;
278 
279  for(int is = 0; is < Ns; ++is){
280 
281  real_t vL[NVCD];
282 
283  if(flag == 0){
284  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
285  vL[ivcd] = 0.0;
286  }
287  }else{
288  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
289  vL[ivcd] = vp[IDX2(Nin5, (ivcd + NVCD * is), site)];
290  }
291  }
292 
293  idir = 1;
294 
295  if ((iy > 0) || (do_comm[idir] == 0)) {
296  int iy2 = (iy - 1 + Ny) % Ny;
297  int nei = ix + Nx * (iy2 + Ny * izt);
298  real_t bc2 = 1.0;
299  if(iy == 0) bc2 = bc[1];
300 
301  real_t ut[NDF];
302  load_u(ut, up, nei + Nst_pad * idir);
303 
304  real_t wt[NVCD];
305  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
306  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
307  }
308  mult_wilson_ymb(vL, ut, wt);
309  }
310 
311  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
312  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
313  }
314 
315  } // is loop
316 
317  } // site loop
318  } // acc parallel
319  } // acc data
320 
321 }
322 
323 
324 //====================================================================
325 void mult_domainwall_5din_zpb(
326  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
327  int Ns, int *bc, int *Nsize, int *do_comm, int flag)
328 {
329  int Nx = Nsize[0];
330  int Ny = Nsize[1];
331  int Nz = Nsize[2];
332  int Nt = Nsize[3];
333  int Nst = Nx * Ny * Nz * Nt;
334  int Nst_pad = CEIL_NWP(Nst);
335 
336  int Nin5 = NVCD * Ns;
337 
338  int size = Nin5 * Nst_pad;
339  int size_u = NDF * Nst_pad * NDIM;
340 
341 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
342  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
343  bc[0:NDIM], do_comm[0:NDIM])
344  {
345 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
346  {
347 
348 #pragma acc loop gang worker vector
349  for (int site = 0; site < Nst; ++site) {
350  int Nxy = Nx * Ny;
351 
352  int ixy = site % Nxy;
353  int izt = site / Nxy;
354  int iz = izt % Nz;
355  int it = izt / Nz;
356 
357  int idir;
358 
359  for(int is = 0; is < Ns; ++is){
360 
361  real_t vL[NVCD];
362 
363  if(flag == 0){
364  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
365  vL[ivcd] = 0.0;
366  }
367  }else{
368  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
369  vL[ivcd] = vp[IDX2(Nin5, (ivcd + NVCD * is), site)];
370  }
371  }
372 
373  idir = 2;
374 
375  if ((iz < Nz-1) || (do_comm[idir] == 0)) {
376  int iz2 = (iz + 1) % Nz;
377  int nei = ixy + Nxy * (iz2 + Nz * it);
378  real_t bc2 = 1.0;
379  if(iz == Nz-1) bc2 = bc[2];
380 
381  real_t ut[NDF];
382  load_u(ut, up, site + Nst_pad * idir);
383 
384  real_t wt[NVCD];
385  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
386  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
387  }
388  mult_wilson_zpb(vL, ut, wt);
389  }
390 
391  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
392  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
393  }
394 
395  } // is loop
396 
397  } // site loop
398  } // acc parallel
399  } // acc data
400 
401 }
402 
403 
404 //====================================================================
405 void mult_domainwall_5din_zmb(
406  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
407  int Ns, int *bc, int *Nsize, int *do_comm, int flag)
408 {
409  int Nx = Nsize[0];
410  int Ny = Nsize[1];
411  int Nz = Nsize[2];
412  int Nt = Nsize[3];
413  int Nst = Nx * Ny * Nz * Nt;
414  int Nst_pad = CEIL_NWP(Nst);
415 
416  int Nin5 = NVCD * Ns;
417 
418  int size = Nin5 * Nst_pad;
419  int size_u = NDF * Nst_pad * NDIM;
420 
421 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
422  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
423  bc[0:NDIM], do_comm[0:NDIM])
424  {
425 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
426  {
427 
428 #pragma acc loop gang worker vector
429  for (int site = 0; site < Nst; ++site) {
430  int Nxy = Nx * Ny;
431 
432  int ixy = site % Nxy;
433  int izt = site / Nxy;
434  int iz = izt % Nz;
435  int it = izt / Nz;
436 
437  int idir;
438 
439  for(int is = 0; is < Ns; ++is){
440 
441  real_t vL[NVCD];
442 
443  if(flag == 0){
444  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
445  vL[ivcd] = 0.0;
446  }
447  }else{
448  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
449  vL[ivcd] = vp[IDX2(Nin5, (ivcd + NVCD * is), site)];
450  }
451  }
452 
453  idir = 2;
454 
455  if ((iz > 0) || (do_comm[idir] == 0)) {
456  int iz2 = (iz - 1 + Nz) % Nz;
457  int nei = ixy + Nxy * (iz2 + Nz * it);
458  real_t bc2 = 1.0;
459  if(iz == 0) bc2 = bc[2];
460 
461  real_t ut[NDF];
462  load_u(ut, up, nei + Nst_pad * idir);
463 
464  real_t wt[NVCD];
465  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
466  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
467  }
468  mult_wilson_zmb(vL, ut, wt);
469  }
470 
471  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
472  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
473  }
474 
475  } // is loop
476 
477  } // site loop
478  } // acc parallel
479  } // acc data
480 
481 }
482 
483 
484 //====================================================================
485 void mult_domainwall_5din_tpb_dirac(
486  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
487  int Ns, int *bc, int *Nsize, int *do_comm, int flag)
488 {
489  int Nx = Nsize[0];
490  int Ny = Nsize[1];
491  int Nz = Nsize[2];
492  int Nt = Nsize[3];
493  int Nst = Nx * Ny * Nz * Nt;
494  int Nst_pad = CEIL_NWP(Nst);
495 
496  int Nin5 = NVCD * Ns;
497 
498  int size = Nin5 * Nst_pad;
499  int size_u = NDF * Nst_pad * NDIM;
500 
501 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
502  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
503  bc[0:NDIM], do_comm[0:NDIM])
504  {
505 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
506  {
507 
508 #pragma acc loop gang worker vector
509  for (int site = 0; site < Nst; ++site) {
510  int Nxy = Nx * Ny;
511  int Nxyz = Nxy * Nz;
512 
513  int izt = site / Nxy;
514  int it = izt / Nz;
515  int ixyz = site % Nxyz;
516 
517  int idir;
518 
519  for(int is = 0; is < Ns; ++is){
520 
521  real_t vL[NVCD];
522 
523  if(flag == 0){
524  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
525  vL[ivcd] = 0.0;
526  }
527  }else{
528  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
529  vL[ivcd] = vp[IDX2(Nin5, (ivcd + NVCD * is), site)];
530  }
531  }
532 
533  idir = 3;
534 
535  if ((it < Nt-1) || (do_comm[idir] == 0)) {
536  int it2 = (it + 1) % Nt;
537  int nei = ixyz + Nxyz * it2;
538  real_t bc2 = 1.0;
539  if(it == Nt-1) bc2 = bc[3];
540 
541  real_t ut[NDF];
542  load_u(ut, up, site + Nst_pad * idir);
543 
544  real_t wt[NVCD];
545  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
546  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
547  }
548  mult_wilson_tpb_dirac(vL, ut, wt);
549  }
550 
551  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
552  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
553  }
554 
555  } // is loop
556 
557  } // site loop
558  } // acc parallel
559  } // acc data
560 
561 }
562 
563 
564 //====================================================================
565 void mult_domainwall_5din_tmb_dirac(
566  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
567  int Ns, int *bc, int *Nsize, int *do_comm, int flag)
568 {
569  int Nx = Nsize[0];
570  int Ny = Nsize[1];
571  int Nz = Nsize[2];
572  int Nt = Nsize[3];
573  int Nst = Nx * Ny * Nz * Nt;
574  int Nst_pad = CEIL_NWP(Nst);
575 
576  int Nin5 = NVCD * Ns;
577 
578  int size = Nin5 * Nst_pad;
579  int size_u = NDF * Nst_pad * NDIM;
580 
581 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
582  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
583  bc[0:NDIM], do_comm[0:NDIM])
584  {
585 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
586  {
587 
588 #pragma acc loop gang worker vector
589  for (int site = 0; site < Nst; ++site) {
590  int Nxy = Nx * Ny;
591  int Nxyz = Nxy * Nz;
592 
593  int izt = site / Nxy;
594  int it = izt / Nz;
595  int ixyz = site % Nxyz;
596 
597  int idir;
598 
599  for(int is = 0; is < Ns; ++is){
600 
601  real_t vL[NVCD];
602 
603  if(flag == 0){
604  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
605  vL[ivcd] = 0.0;
606  }
607  }else{
608  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
609  vL[ivcd] = vp[IDX2(Nin5, (ivcd + NVCD * is), site)];
610  }
611  }
612 
613  idir = 3;
614 
615  if ((it > 0) || (do_comm[idir] == 0)) {
616  int it2 = (it - 1 + Nt) % Nt;
617  int nei = ixyz + Nxyz * it2;
618  real_t bc2 = 1.0;
619  if(it == 0) bc2 = bc[3];
620 
621  real_t ut[NDF];
622  load_u(ut, up, nei + Nst_pad * idir);
623 
624  real_t wt[NVCD];
625  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
626  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
627  }
628  mult_wilson_tmb_dirac(vL, ut, wt);
629  }
630 
631  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
632  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
633  }
634 
635  } // is loop
636 
637  } // site loop
638  } // acc parallel
639  } // acc data
640 
641 }
642 
643 
644 //====================================================================
645 void mult_domainwall_5din_xp1(real_t *RESTRICT buf_xp,
646  real_t *RESTRICT up, real_t *RESTRICT wp,
647  int Ns, int *bc, int *Nsize)
648 {
649  int Nx = Nsize[0];
650  int Ny = Nsize[1];
651  int Nz = Nsize[2];
652  int Nt = Nsize[3];
653  int Nst = Nx * Ny * Nz * Nt;
654  int Nst_pad = CEIL_NWP(Nst);
655 
656  int Nin5 = NVCD * Ns;
657  int Nin5bd = NVC * ND2 * Ns;
658 
659  int size = Nin5 * Nst_pad;
660  int size_u = NDF * Nst_pad * NDIM;
661 
662  int size_bx = Nin5bd * CEIL_NWP(Ny * Nz * Nt);
663 
664 #pragma acc data present(up[0:size_u], wp[0:size], \
665  buf_xp[0:size_bx]) \
666  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
667  {
668 
669  int Nyzt = Ny * Nz * Nt;
670 
671 #pragma acc parallel \
672  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
673 #pragma acc loop gang worker vector
674  for (int iyzt = 0; iyzt < Nyzt; ++iyzt) {
675  int ix = 0;
676  int site = ix + Nx * iyzt;
677  real_t bc2 = bc[0];
678  for (int is = 0; is < Ns; ++is) {
679  real_t wt[NVCD], vt[NVC * ND2];
680  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
681  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
682  }
683  mult_wilson_xp1(vt, wt);
684  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
685  buf_xp[IDX2(Nin5bd, (ivcd + NVCD2 * is), iyzt)] = bc2 * vt[ivcd];
686  }
687  }
688  }
689 
690  }
691 }
692 
693 //====================================================================
694 void mult_domainwall_5din_xm1(real_t *RESTRICT buf_xm,
695  real_t *RESTRICT up, real_t *RESTRICT wp,
696  int Ns, int *bc, int *Nsize)
697 {
698  int Nx = Nsize[0];
699  int Ny = Nsize[1];
700  int Nz = Nsize[2];
701  int Nt = Nsize[3];
702  int Nst = Nx * Ny * Nz * Nt;
703  int Nst_pad = CEIL_NWP(Nst);
704 
705  int Nin5 = NVCD * Ns;
706  int Nin5bd = NVC * ND2 * Ns;
707 
708  int size = Nin5 * Nst_pad;
709  int size_u = NDF * Nst_pad * NDIM;
710  int size_bx = Nin5bd * CEIL_NWP(Ny * Nz * Nt);
711 
712 #pragma acc data present(up[0:size_u], wp[0:size], \
713  buf_xm[0:size_bx]) \
714  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
715  {
716 
717  int idir = 0;
718  int Nyzt = Ny * Nz * Nt;
719 
720 #pragma acc parallel \
721  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
722 #pragma acc loop gang worker vector
723  for (int iyzt = 0; iyzt < Nyzt; ++iyzt) {
724  int ix = Nx-1;
725  int site = ix + Nx * iyzt;
726  real_t ut[NDF];
727  load_u(ut, up, site + Nst_pad * idir);
728  for (int is = 0; is < Ns; ++is) {
729  real_t wt[NVCD], vt[NVC * ND2];
730  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
731  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
732  }
733  mult_wilson_xm1(vt, ut, wt);
734  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
735  buf_xm[IDX2(Nin5bd, (ivcd + NVCD2 * is), iyzt)] = vt[ivcd];
736  }
737  }
738  }
739  }
740 
741 }
742 
743 //====================================================================
744 void mult_domainwall_5din_yp1(real_t *RESTRICT buf_yp,
745  real_t *RESTRICT up, real_t *RESTRICT wp,
746  int Ns, int *bc, int *Nsize)
747 {
748  int Nx = Nsize[0];
749  int Ny = Nsize[1];
750  int Nz = Nsize[2];
751  int Nt = Nsize[3];
752  int Nst = Nx * Ny * Nz * Nt;
753  int Nst_pad = CEIL_NWP(Nst);
754 
755  int Nin5 = NVCD * Ns;
756  int Nin5bd = NVC * ND2 * Ns;
757 
758  int size = Nin5 * Nst_pad;
759  int size_u = NDF * Nst_pad * NDIM;
760  int size_by = Nin5bd * CEIL_NWP(Nx * Nz * Nt);
761 
762 #pragma acc data present(up[0:size_u], wp[0:size], \
763  buf_yp[0:size_by]) \
764  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
765  {
766 
767  int Nxzt = Nx * Nz * Nt;
768 
769 #pragma acc parallel \
770  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
771 #pragma acc loop gang worker vector
772  for (int ixzt = 0; ixzt < Nxzt; ++ixzt) {
773  int iy = 0;
774  int ix = ixzt % Nx;
775  int izt = ixzt / Nx;
776  int site = ix + Nx * (iy + Ny * izt);
777  real_t bc2 = bc[1];
778  for (int is = 0; is < Ns; ++is) {
779  real_t wt[NVCD], vt[NVC * ND2];
780  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
781  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
782  }
783  mult_wilson_yp1(vt, wt);
784  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
785  buf_yp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixzt)] = bc2 * vt[ivcd];
786  }
787  }
788  }
789  }
790 }
791 
792 //====================================================================
793 void mult_domainwall_5din_ym1(real_t *RESTRICT buf_ym,
794  real_t *RESTRICT up, real_t *RESTRICT wp,
795  int Ns, int *bc, int *Nsize)
796 {
797  int Nx = Nsize[0];
798  int Ny = Nsize[1];
799  int Nz = Nsize[2];
800  int Nt = Nsize[3];
801  int Nst = Nx * Ny * Nz * Nt;
802  int Nst_pad = CEIL_NWP(Nst);
803 
804  int Nin5 = NVCD * Ns;
805  int Nin5bd = NVC * ND2 * Ns;
806 
807  int size = Nin5 * Nst_pad;
808  int size_u = NDF * Nst_pad * NDIM;
809  int size_by = Nin5bd * CEIL_NWP(Nx * Nz * Nt);
810 
811 #pragma acc data present(up[0:size_u], wp[0:size], \
812  buf_ym[0:size_by]) \
813  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
814  {
815 
816  int idir = 1;
817  int Nxzt = Nx * Nz * Nt;
818 
819 #pragma acc parallel \
820  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
821 #pragma acc loop gang worker vector
822  for (int ixzt = 0; ixzt < Nxzt; ++ixzt) {
823  int iy = Ny-1;
824  int ix = ixzt % Nx;
825  int izt = ixzt / Nx;
826  int site = ix + Nx * (iy + Ny * izt);
827  real_t ut[NDF];
828  load_u(ut, up, site + Nst_pad * idir);
829  for (int is = 0; is < Ns; ++is) {
830  real_t wt[NVCD], vt[NVC * ND2];
831  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
832  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
833  }
834  mult_wilson_ym1(vt, ut, wt);
835  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
836  buf_ym[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixzt)] = vt[ivcd];
837  }
838  }
839  }
840  }
841 
842 }
843 
844 //====================================================================
845 void mult_domainwall_5din_zp1(real_t *RESTRICT buf_zp,
846  real_t *RESTRICT up, real_t *RESTRICT wp,
847  int Ns, int *bc, int *Nsize)
848 {
849  int Nx = Nsize[0];
850  int Ny = Nsize[1];
851  int Nz = Nsize[2];
852  int Nt = Nsize[3];
853  int Nst = Nx * Ny * Nz * Nt;
854  int Nst_pad = CEIL_NWP(Nst);
855 
856  int Nin5 = NVCD * Ns;
857  int Nin5bd = NVC * ND2 * Ns;
858 
859  int size = Nin5 * Nst_pad;
860  int size_u = NDF * Nst_pad * NDIM;
861  int size_bz = Nin5bd * CEIL_NWP(Nx * Ny * Nt);
862 
863 #pragma acc data present(up[0:size_u], wp[0:size], buf_zp[0:size_bz])\
864  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
865  {
866  int Nxy = Nx * Ny;
867  int Nxyt = Nx * Ny * Nt;
868 
869 #pragma acc parallel \
870  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
871 #pragma acc loop gang worker vector
872  for (int ixyt = 0; ixyt < Nxyt; ++ixyt) {
873  int iz = 0;
874  int ixy = ixyt % Nxy;
875  int it = ixyt / Nxy;
876  int site = ixy + Nxy * (iz + Nz * it);
877  real_t bc2 = bc[2];
878  for (int is = 0; is < Ns; ++is) {
879  real_t wt[NVCD], vt[NVC * ND2];
880  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
881  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
882  }
883  mult_wilson_zp1(vt, wt);
884  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
885  buf_zp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyt)] = bc2 * vt[ivcd];
886  }
887  }
888  }
889  }
890 
891 }
892 
893 //====================================================================
894 void mult_domainwall_5din_zm1(real_t *RESTRICT buf_zm,
895  real_t *RESTRICT up, real_t *RESTRICT wp,
896  int Ns, int *bc, int *Nsize)
897 {
898  int Nx = Nsize[0];
899  int Ny = Nsize[1];
900  int Nz = Nsize[2];
901  int Nt = Nsize[3];
902  int Nst = Nx * Ny * Nz * Nt;
903  int Nst_pad = CEIL_NWP(Nst);
904 
905  int Nin5 = NVCD * Ns;
906  int Nin5bd = NVC * ND2 * Ns;
907 
908  int size = Nin5 * Nst_pad;
909  int size_u = NDF * Nst_pad * NDIM;
910  int size_bz = Nin5bd * CEIL_NWP(Nx * Ny * Nt);
911 
912 #pragma acc data present(up[0:size_u], wp[0:size], \
913  buf_zm[0:size_bz]) \
914  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
915  {
916 
917  int idir = 2;
918  int Nxy = Nx * Ny;
919  int Nxyt = Nx * Ny * Nt;
920 
921 #pragma acc parallel \
922  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
923 #pragma acc loop gang worker vector
924  for (int ixyt = 0; ixyt < Nxyt; ++ixyt) {
925  int iz = Nz-1;
926  int ixy = ixyt % Nxy;
927  int it = ixyt / Nxy;
928  int site = ixy + Nxy * (iz + Nz * it);
929  real_t ut[NDF];
930  load_u(ut, up, site + Nst_pad * idir);
931  for (int is = 0; is < Ns; ++is) {
932  real_t wt[NVCD], vt[NVC * ND2];
933  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
934  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
935  }
936  mult_wilson_zm1(vt, ut, wt);
937  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
938  buf_zm[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyt)] = vt[ivcd];
939  }
940  }
941  }
942  }
943 
944 }
945 
946 //====================================================================
947 void mult_domainwall_5din_tp1_dirac(
948  real_t *RESTRICT buf_tp,
949  real_t *RESTRICT up, real_t *RESTRICT wp,
950  int Ns, int *bc, int *Nsize)
951 {
952  int Nx = Nsize[0];
953  int Ny = Nsize[1];
954  int Nz = Nsize[2];
955  int Nt = Nsize[3];
956  int Nst = Nx * Ny * Nz * Nt;
957  int Nst_pad = CEIL_NWP(Nst);
958 
959  int Nin5 = NVCD * Ns;
960  int Nin5bd = NVC * ND2 * Ns;
961 
962  int size = Nin5 * Nst_pad;
963  int size_u = NDF * Nst_pad * NDIM;
964  int size_bt = Nin5bd * CEIL_NWP(Nx * Ny * Nz);
965 
966 #pragma acc data present(up[0:size_u], wp[0:size], \
967  buf_tp[0:size_bt] ) \
968  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
969  {
970  int Nxyz = Nx * Ny * Nz;
971 
972 #pragma acc parallel \
973  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
974 #pragma acc loop gang worker vector
975  for (int ixyz = 0; ixyz < Nxyz; ++ixyz) {
976  int it = 0;
977  int site = ixyz + Nxyz * it;
978  real_t bc2 = bc[3];
979  for (int is = 0; is < Ns; ++is) {
980  real_t wt[NVCD], vt[NVC * ND2];
981  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
982  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
983  }
984  mult_wilson_tp1_dirac(vt, wt);
985  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
986  buf_tp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyz)] = bc2 * vt[ivcd];
987  }
988  }
989  }
990  }
991 
992 }
993 
994 //====================================================================
995 void mult_domainwall_5din_tm1_dirac(
996  real_t *RESTRICT buf_tm,
997  real_t *RESTRICT up, real_t *RESTRICT wp,
998  int Ns, int *bc, int *Nsize)
999 {
1000  int Nx = Nsize[0];
1001  int Ny = Nsize[1];
1002  int Nz = Nsize[2];
1003  int Nt = Nsize[3];
1004  int Nst = Nx * Ny * Nz * Nt;
1005  int Nst_pad = CEIL_NWP(Nst);
1006 
1007  int Nin5 = NVCD * Ns;
1008  int Nin5bd = NVC * ND2 * Ns;
1009 
1010  int size = Nin5 * Nst_pad;
1011  int size_u = NDF * Nst_pad * NDIM;
1012  int size_bt = Nin5bd * CEIL_NWP(Nx * Ny * Nz);
1013 
1014 #pragma acc data present(up[0:size_u], wp[0:size], \
1015  buf_tm[0:size_bt] ) \
1016  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1017  {
1018 
1019  int idir = 3;
1020  int Nxyz = Nx * Ny * Nz;
1021 
1022 #pragma acc parallel \
1023  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1024 #pragma acc loop gang worker vector
1025  for (int ixyz = 0; ixyz < Nxyz; ++ixyz) {
1026  int it = Nt-1;
1027  int site = ixyz + Nxyz * it;
1028  real_t ut[NDF];
1029  load_u(ut, up, site + Nst_pad * idir);
1030  for (int is = 0; is < Ns; ++is) {
1031  real_t wt[NVCD], vt[NVC * ND2];
1032  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1033  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
1034  }
1035  mult_wilson_tm1_dirac(vt, ut, wt);
1036  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
1037  buf_tm[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyz)] = vt[ivcd];
1038  }
1039  }
1040  }
1041  }
1042 
1043 }
1044 
1045 //====================================================================
1046 void mult_domainwall_5din_xp2(real_t *RESTRICT vp, real_t *RESTRICT up,
1047  real_t *RESTRICT buf_xp,
1048  int Ns, int *bc, int *Nsize)
1049 {
1050  int Nx = Nsize[0];
1051  int Ny = Nsize[1];
1052  int Nz = Nsize[2];
1053  int Nt = Nsize[3];
1054  int Nst = Nx * Ny * Nz * Nt;
1055  int Nst_pad = CEIL_NWP(Nst);
1056 
1057  int Nin5 = NVCD * Ns;
1058  int Nin5bd = (NVCD/2) * Ns;
1059 
1060  int size = Nin5 * Nst_pad;
1061  int size_u = NDF * Nst_pad * 4;
1062  int size_bx = Nin5bd * CEIL_NWP(Ny * Nz * Nt);
1063 
1064 #pragma acc data present(buf_xp[0:size_bx], \
1065  vp[0:size], up[0:size_u]) \
1066  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1067  {
1068 
1069 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1070  {
1071 
1072 #pragma acc loop gang worker vector
1073  for (int site = 0; site < Nst; ++site) {
1074  int ix = site % Nx;
1075  int iyzt = site / Nx;
1076 
1077  int idir;
1078 
1079  for(int is = 0; is < Ns; ++is){
1080 
1081  real_t vL[NVCD];
1082 
1083  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1084  vL[ivcd] = 0.0;
1085  }
1086 
1087  idir = 0;
1088 
1089  if (ix == Nx-1) {
1090  real_t ut[NDF];
1091  load_u(ut, up, site + Nst_pad * idir);
1092  real_t wt[NVC * ND2];
1093  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1094  wt[ivcd] = buf_xp[IDX2(Nin5bd, (ivcd + NVCD2 * is), iyzt)];
1095  }
1096  mult_wilson_xp2(vL, ut, wt);
1097 
1098  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1099  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] += vL[ivcd];
1100  }
1101  }
1102 
1103  } // is loop
1104  } // site loop
1105 
1106  } // acc parallel
1107  } // acc data
1108 
1109 }
1110 //====================================================================
1111 void mult_domainwall_5din_xm2(real_t *RESTRICT vp, real_t *RESTRICT up,
1112  real_t *RESTRICT buf_xm,
1113  int Ns, int *bc, int *Nsize)
1114 {
1115  int Nx = Nsize[0];
1116  int Ny = Nsize[1];
1117  int Nz = Nsize[2];
1118  int Nt = Nsize[3];
1119  int Nst = Nx * Ny * Nz * Nt;
1120  int Nst_pad = CEIL_NWP(Nst);
1121 
1122  int Nin5 = NVCD * Ns;
1123  int Nin5bd = (NVCD/2) * Ns;
1124 
1125  int size = Nin5 * Nst_pad;
1126  int size_u = NDF * Nst_pad * 4;
1127  int size_bx = Nin5bd * CEIL_NWP(Ny * Nz * Nt);
1128 
1129 #pragma acc data present(buf_xm[0:size_bx], \
1130  vp[0:size], up[0:size_u]) \
1131  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1132  {
1133 
1134 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1135  {
1136 
1137 #pragma acc loop gang worker vector
1138  for (int site = 0; site < Nst; ++site) {
1139  int ix = site % Nx;
1140  int iyzt = site / Nx;
1141 
1142  for(int is = 0; is < Ns; ++is){
1143 
1144  real_t vL[NVCD];
1145 
1146  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1147  vL[ivcd] = 0.0;
1148  }
1149 
1150  if (ix == 0) {
1151  real_t bc2 = bc[0];
1152  real_t wt[NVC * ND2];
1153  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1154  wt[ivcd] = bc2 * buf_xm[IDX2(Nin5bd, (ivcd + NVCD2 * is), iyzt)];
1155  }
1156  mult_wilson_xm2(vL, wt);
1157 
1158  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1159  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] += vL[ivcd];
1160  }
1161  }
1162 
1163  } // is loop
1164  } // site loop
1165 
1166  } // acc parallel
1167  } // acc data
1168 
1169 }
1170 //====================================================================
1171 void mult_domainwall_5din_yp2(real_t *RESTRICT vp, real_t *RESTRICT up,
1172  real_t *RESTRICT buf_yp,
1173  int Ns, int *bc, int *Nsize)
1174 {
1175  int Nx = Nsize[0];
1176  int Ny = Nsize[1];
1177  int Nz = Nsize[2];
1178  int Nt = Nsize[3];
1179  int Nst = Nx * Ny * Nz * Nt;
1180  int Nst_pad = CEIL_NWP(Nst);
1181 
1182  int Nin5 = NVCD * Ns;
1183  int Nin5bd = (NVCD/2) * Ns;
1184 
1185  int size = Nin5 * Nst_pad;
1186  int size_u = NDF * Nst_pad * 4;
1187  int size_by = Nin5bd * CEIL_NWP(Nx * Nz * Nt);
1188 
1189 #pragma acc data present(buf_yp[0:size_by], \
1190  vp[0:size], up[0:size_u]) \
1191  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1192  {
1193 
1194 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1195  {
1196  int Nxy = Nx * Ny;
1197 
1198 #pragma acc loop gang worker vector
1199  for (int site = 0; site < Nst; ++site) {
1200  int ix = site % Nx;
1201  int iyzt = site / Nx;
1202  int iy = iyzt % Ny;
1203  int izt = site / Nxy;
1204 
1205  int idir;
1206 
1207  for(int is = 0; is < Ns; ++is){
1208 
1209  real_t vL[NVCD];
1210 
1211  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1212  vL[ivcd] = 0.0;
1213  }
1214 
1215  idir = 1;
1216  int ixzt = ix + Nx * izt;
1217 
1218  if (iy == Ny-1) {
1219  real_t ut[NDF];
1220  load_u(ut, up, site + Nst_pad * idir);
1221  real_t wt[NVC * ND2];
1222  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1223  wt[ivcd] = buf_yp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixzt)];
1224  }
1225  mult_wilson_yp2(vL, ut, wt);
1226 
1227  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1228  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] += vL[ivcd];
1229  }
1230  }
1231 
1232  } // is loop
1233  } // site loop
1234 
1235  } // acc parallel
1236  } // acc data
1237 
1238 }
1239 //====================================================================
1240 void mult_domainwall_5din_ym2(real_t *RESTRICT vp, real_t *RESTRICT up,
1241  real_t *RESTRICT buf_ym,
1242  int Ns, int *bc, int *Nsize)
1243 {
1244  int Nx = Nsize[0];
1245  int Ny = Nsize[1];
1246  int Nz = Nsize[2];
1247  int Nt = Nsize[3];
1248  int Nst = Nx * Ny * Nz * Nt;
1249  int Nst_pad = CEIL_NWP(Nst);
1250 
1251  int Nin5 = NVCD * Ns;
1252  int Nin5bd = (NVCD/2) * Ns;
1253 
1254  int size = Nin5 * Nst_pad;
1255  int size_u = NDF * Nst_pad * 4;
1256  int size_by = Nin5bd * CEIL_NWP(Nx * Nz * Nt);
1257 
1258 #pragma acc data present(buf_ym[0:size_by], \
1259  vp[0:size], up[0:size_u]) \
1260  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1261  {
1262 
1263 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1264  {
1265  int Nxy = Nx * Ny;
1266 
1267 #pragma acc loop gang worker vector
1268  for (int site = 0; site < Nst; ++site) {
1269  int ix = site % Nx;
1270  int iyzt = site / Nx;
1271  int iy = iyzt % Ny;
1272  int izt = site / Nxy;
1273 
1274  for(int is = 0; is < Ns; ++is){
1275 
1276  real_t vL[NVCD];
1277 
1278  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1279  vL[ivcd] = 0.0;
1280  }
1281 
1282  int ixzt = ix + Nx * izt;
1283 
1284  if (iy == 0) {
1285  real_t bc2 = bc[1];
1286  real_t wt[NVC * ND2];
1287  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1288  wt[ivcd] = bc2 * buf_ym[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixzt)];
1289  }
1290  mult_wilson_ym2(vL, wt);
1291 
1292  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1293  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] += vL[ivcd];
1294  }
1295  }
1296 
1297  } // is loop
1298  } // site loop
1299 
1300  } // acc parallel
1301  } // acc data
1302 
1303 }
1304 
1305 //====================================================================
1306 void mult_domainwall_5din_zp2(real_t *RESTRICT vp, real_t *RESTRICT up,
1307  real_t *RESTRICT buf_zp,
1308  int Ns, int *bc, int *Nsize)
1309 {
1310  int Nx = Nsize[0];
1311  int Ny = Nsize[1];
1312  int Nz = Nsize[2];
1313  int Nt = Nsize[3];
1314  int Nst = Nx * Ny * Nz * Nt;
1315  int Nst_pad = CEIL_NWP(Nst);
1316 
1317  int Nin5 = NVCD * Ns;
1318  int Nin5bd = (NVCD/2) * Ns;
1319 
1320  int size = Nin5 * Nst_pad;
1321  int size_u = NDF * Nst_pad * 4;
1322  int size_bz = Nin5bd * CEIL_NWP(Nx * Ny * Nt);
1323 
1324 #pragma acc data present(buf_zp[0:size_bz], \
1325  vp[0:size], up[0:size_u]) \
1326  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1327  {
1328 
1329 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1330  {
1331  int Nxy = Nx * Ny;
1332 
1333 #pragma acc loop gang worker vector
1334  for (int site = 0; site < Nst; ++site) {
1335  int ixy = site % Nxy;
1336  int izt = site / Nxy;
1337  int iz = izt % Nz;
1338  int it = izt / Nz;
1339 
1340  for(int is = 0; is < Ns; ++is){
1341 
1342  real_t vL[NVCD];
1343 
1344  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1345  vL[ivcd] = 0.0;
1346  }
1347 
1348  int idir = 2;
1349  int ixyt = ixy + Nxy * it;
1350  if (iz == Nz-1) {
1351  real_t ut[NDF];
1352  load_u(ut, up, site + Nst_pad * idir);
1353  real_t wt[NVC * ND2];
1354  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1355  wt[ivcd] = buf_zp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyt)];
1356  }
1357  mult_wilson_zp2(vL, ut, wt);
1358 
1359  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1360  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] += vL[ivcd];
1361  }
1362  }
1363 
1364  } // is loop
1365  } // site loop
1366 
1367  } // acc parallel
1368  } // acc data
1369 
1370 }
1371 //====================================================================
1372 void mult_domainwall_5din_zm2(real_t *RESTRICT vp, real_t *RESTRICT up,
1373  real_t *RESTRICT buf_zm,
1374  int Ns, int *bc, int *Nsize)
1375 {
1376  int Nx = Nsize[0];
1377  int Ny = Nsize[1];
1378  int Nz = Nsize[2];
1379  int Nt = Nsize[3];
1380  int Nst = Nx * Ny * Nz * Nt;
1381  int Nst_pad = CEIL_NWP(Nst);
1382 
1383  int Nin5 = NVCD * Ns;
1384  int Nin5bd = (NVCD/2) * Ns;
1385 
1386  int size = Nin5 * Nst_pad;
1387  int size_u = NDF * Nst_pad * 4;
1388  int size_bz = Nin5bd * CEIL_NWP(Nx * Ny * Nt);
1389 
1390 #pragma acc data present(buf_zm[0:size_bz], \
1391  vp[0:size], up[0:size_u]) \
1392  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1393  {
1394 
1395 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1396  {
1397  int Nxy = Nx * Ny;
1398 
1399 #pragma acc loop gang worker vector
1400  for (int site = 0; site < Nst; ++site) {
1401  int ixy = site % Nxy;
1402  int izt = site / Nxy;
1403  int iz = izt % Nz;
1404  int it = izt / Nz;
1405 
1406  for(int is = 0; is < Ns; ++is){
1407 
1408  real_t vL[NVCD];
1409 
1410  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1411  vL[ivcd] = 0.0;
1412  }
1413 
1414  int ixyt = ixy + Nxy * it;
1415 
1416  if (iz == 0) {
1417  real_t bc2 = bc[2];
1418  real_t wt[NVC * ND2];
1419  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1420  wt[ivcd] = bc2 * buf_zm[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyt)];
1421  }
1422  mult_wilson_zm2(vL, wt);
1423 
1424  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1425  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] += vL[ivcd];
1426  }
1427  }
1428 
1429  } // is loop
1430  } // site loop
1431 
1432  } // acc parallel
1433  } // acc data
1434 
1435 }
1436 //====================================================================
1437 void mult_domainwall_5din_tp2_dirac(
1438  real_t *RESTRICT vp, real_t *RESTRICT up,
1439  real_t *RESTRICT buf_tp,
1440  int Ns, int *bc, int *Nsize)
1441 {
1442  int Nx = Nsize[0];
1443  int Ny = Nsize[1];
1444  int Nz = Nsize[2];
1445  int Nt = Nsize[3];
1446  int Nst = Nx * Ny * Nz * Nt;
1447  int Nst_pad = CEIL_NWP(Nst);
1448 
1449  int Nin5 = NVCD * Ns;
1450  int Nin5bd = (NVCD/2) * Ns;
1451 
1452  int size = Nin5 * Nst_pad;
1453  int size_u = NDF * Nst_pad * 4;
1454  int size_bt = Nin5bd * CEIL_NWP(Nx * Ny * Nz);
1455 
1456 #pragma acc data present(buf_tp[0:size_bt], \
1457  vp[0:size], up[0:size_u]) \
1458  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1459  {
1460 
1461 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1462  {
1463  int Nxy = Nx * Ny;
1464  int Nxyz = Nx * Ny * Nz;
1465 
1466 #pragma acc loop gang worker vector
1467  for (int site = 0; site < Nst; ++site) {
1468  int izt = site / Nxy;
1469  int it = izt / Nz;
1470  int ixyz = site % Nxyz;
1471 
1472  for(int is = 0; is < Ns; ++is){
1473 
1474  real_t vL[NVCD];
1475 
1476  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1477  vL[ivcd] = 0.0;
1478  }
1479 
1480  int idir = 3;
1481  if (it == Nt-1) {
1482  real_t ut[NDF];
1483  load_u(ut, up, site + Nst_pad * idir);
1484  real_t wt[NVC * ND2];
1485  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1486  wt[ivcd] = buf_tp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyz)];
1487  }
1488  mult_wilson_tp2_dirac(vL, ut, wt);
1489 
1490  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1491  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] += vL[ivcd];
1492  }
1493  }
1494 
1495  } // is loop
1496  } // site loop
1497 
1498  } // acc parallel
1499  } // acc data
1500 
1501 }
1502 //====================================================================
1503 void mult_domainwall_5din_tm2_dirac(
1504  real_t *RESTRICT vp, real_t *RESTRICT up,
1505  real_t *RESTRICT buf_tm,
1506  int Ns, int *bc, int *Nsize)
1507 {
1508  int Nx = Nsize[0];
1509  int Ny = Nsize[1];
1510  int Nz = Nsize[2];
1511  int Nt = Nsize[3];
1512  int Nst = Nx * Ny * Nz * Nt;
1513  int Nst_pad = CEIL_NWP(Nst);
1514 
1515  int Nin5 = NVCD * Ns;
1516  int Nin5bd = (NVCD/2) * Ns;
1517 
1518  int size = Nin5 * Nst_pad;
1519  int size_u = NDF * Nst_pad * 4;
1520  int size_bt = Nin5bd * CEIL_NWP(Nx * Ny * Nz);
1521 
1522 #pragma acc data present(buf_tm[0:size_bt], \
1523  vp[0:size], up[0:size_u]) \
1524  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, bc[0:NDIM])
1525  {
1526 
1527 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1528  {
1529  int Nxy = Nx * Ny;
1530  int Nxyz = Nx * Ny * Nz;
1531 
1532 #pragma acc loop gang worker vector
1533  for (int site = 0; site < Nst; ++site) {
1534  int izt = site / Nxy;
1535  int it = izt / Nz;
1536  int ixyz = site % Nxyz;
1537 
1538  for(int is = 0; is < Ns; ++is){
1539 
1540  real_t vL[NVCD];
1541 
1542  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1543  vL[ivcd] = 0.0;
1544  }
1545 
1546  if (it == 0) {
1547  real_t bc2 = bc[3];
1548  real_t wt[NVC * ND2];
1549  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1550  wt[ivcd] = bc2 * buf_tm[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyz)];
1551  }
1552  mult_wilson_tm2_dirac(vL, wt);
1553 
1554  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1555  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] += vL[ivcd];
1556  }
1557  }
1558 
1559  } // is loop
1560  } // site loop
1561 
1562  } // acc parallel
1563  } // acc data
1564 
1565 }
1566 
1567 #endif
1568 //============================================================END=====
BridgeACC::mult_wilson_yp2
void mult_wilson_yp2(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:467
mult_domainwall_5din_xp1
void mult_domainwall_5din_xp1(real_t *RESTRICT buf_xp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:645
iy
int iy
Definition: mult_Wilson_xyz_openacc-inc.h:239
ixy
int ixy
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:513
wt
real_t wt[NVCD]
Definition: mult_Clover_csw_chiral_openacc-inc.h:9
BridgeACC::mult_wilson_xpb
void mult_wilson_xpb(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:130
izt
int izt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:262
mult_domainwall_5din_zp2
void mult_domainwall_5din_zp2(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT buf_zp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:1306
mult_domainwall_5din_zm2
void mult_domainwall_5din_zm2(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT buf_zm, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:1372
NVCD
#define NVCD
Definition: define_params_SU3.h:20
mult_domainwall_5din_xpb
void mult_domainwall_5din_xpb(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int flag)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:14
mult_domainwall_5din_ymb
void mult_domainwall_5din_ymb(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int flag)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:245
BridgeACC::mult_wilson_xm2
void mult_wilson_xm2(double *RESTRICT v2, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:282
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
mult_domainwall_5din_yp2
void mult_domainwall_5din_yp2(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT buf_yp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:1171
mult_domainwall_5din_xm2
void mult_domainwall_5din_xm2(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT buf_xm, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:1111
BridgeACC::mult_wilson_yp1
void mult_wilson_yp1(double *RESTRICT buf, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:414
mult_domainwall_5din_xmb
void mult_domainwall_5din_xmb(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int flag)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:90
mult_domainwall_5din_zpb
void mult_domainwall_5din_zpb(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int flag)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:325
mult_domainwall_5din_zp1
void mult_domainwall_5din_zp1(real_t *RESTRICT buf_zp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:845
BridgeACC::mult_wilson_zp1
void mult_wilson_zp1(double *RESTRICT buf, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:836
BridgeACC::mult_wilson_tp2_dirac
void mult_wilson_tp2_dirac(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1311
mult_domainwall_5din_ym2
void mult_domainwall_5din_ym2(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT buf_ym, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:1240
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
BridgeACC::mult_wilson_tm2_dirac
void mult_wilson_tm2_dirac(double *RESTRICT v2, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1522
iyzt
int iyzt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:14
ix
int ix
Definition: mult_Wilson_xyz_openacc-inc.h:16
BridgeACC::mult_wilson_tmb_dirac
void mult_wilson_tmb_dirac(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1569
BridgeACC::mult_wilson_xmb
void mult_wilson_xmb(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:334
idir
idir
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:264
mult_domainwall_5din_tm1_dirac
void mult_domainwall_5din_tm1_dirac(real_t *RESTRICT buf_tm, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:995
BridgeACC::mult_wilson_xm1
void mult_wilson_xm1(double *RESTRICT buf, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:209
BridgeACC::mult_wilson_ym2
void mult_wilson_ym2(double *RESTRICT v2, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:696
bc2
int bc2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:62
BridgeACC::mult_wilson_zpb
void mult_wilson_zpb(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:962
BridgeACC::mult_wilson_tpb_dirac
void mult_wilson_tpb_dirac(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1373
mult_domainwall_5din_tm2_dirac
void mult_domainwall_5din_tm2_dirac(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT buf_tm, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:1503
mult_domainwall_5din_yp1
void mult_domainwall_5din_yp1(real_t *RESTRICT buf_yp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:744
mult_domainwall_5din_xm1
void mult_domainwall_5din_xm1(real_t *RESTRICT buf_xm, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:694
BridgeACC::mult_wilson_tm1_dirac
void mult_wilson_tm1_dirac(double *RESTRICT buf, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1450
it
int it
Definition: mult_Wilson_xyz_openacc-inc.h:461
BridgeACC::mult_wilson_zmb
void mult_wilson_zmb(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1177
ix2
int ix2
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:138
BridgeACC::mult_wilson_tp1_dirac
void mult_wilson_tp1_dirac(double *RESTRICT buf, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1261
BridgeACC::mult_wilson_xp2
void mult_wilson_xp2(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:64
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
BridgeACC::mult_wilson_ypb
void mult_wilson_ypb(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:537
mult_domainwall_5din_tp2_dirac
void mult_domainwall_5din_tp2_dirac(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT buf_tp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:1437
IDX2
#define IDX2(nin, in, ist)
Definition: define_index.h:28
BridgeACC::mult_wilson_zm2
void mult_wilson_zm2(double *RESTRICT v2, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1121
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
BridgeACC::mult_wilson_ymb
void mult_wilson_ymb(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:752
load_u
void load_u(real_t *ut, real_t *up, int site)
Definition: mult_Wilson_inline_openacc-inc.h:26
mult_domainwall_5din_tpb_dirac
void mult_domainwall_5din_tpb_dirac(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int flag)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:485
mult_domainwall_5din_xp2
void mult_domainwall_5din_xp2(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT buf_xp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:1046
BridgeACC::mult_wilson_zp2
void mult_wilson_zp2(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:891
iz
int iz
Definition: mult_Wilson_xyz_openacc-inc.h:462
mult_domainwall_5din_ypb
void mult_domainwall_5din_ypb(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int flag)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:167
mult_domainwall_5din_zm1
void mult_domainwall_5din_zm1(real_t *RESTRICT buf_zm, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:894
BridgeACC::mult_wilson_ym1
void mult_wilson_ym1(double *RESTRICT buf, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:621
NDIM
#define NDIM
Definition: contract_4spinor.cpp:18
ixyz
int ixyz
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:13
mult_domainwall_5din_zmb
void mult_domainwall_5din_zmb(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int flag)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:405
BridgeACC::mult_wilson_zm1
void mult_wilson_zm1(double *RESTRICT buf, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1045
BridgeACC::mult_wilson_xp1
void mult_wilson_xp1(double *RESTRICT buf, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:14
mult_domainwall_5din_tmb_dirac
void mult_domainwall_5din_tmb_dirac(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int flag)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:565
ND2
#define ND2
Definition: define_params_SU3.h:18
mult_domainwall_5din_tp1_dirac
void mult_domainwall_5din_tp1_dirac(real_t *RESTRICT buf_tp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:947
mult_domainwall_5din_ym1
void mult_domainwall_5din_ym1(real_t *RESTRICT buf_ym, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize)
Definition: mult_Domainwall_5din_dir_openacc-inc.h:793