Bridge++  Ver.2.1.3
mult_Wilson_eo_openacc-inc.h
Go to the documentation of this file.
1 
10 #define MULT_GXr(u0,u1,u2,u3,u4,u5,v0,v1,v2,v3,v4,v5) (u0*v0-u1*v1 + u2*v2-u3*v3 + u4*v4-u5*v5)
11 #define MULT_GXi(u0,u1,u2,u3,u4,u5,v0,v1,v2,v3,v4,v5) (u0*v1+u1*v0 + u2*v3+u3*v2 + u4*v5+u5*v4)
12 #define MULT_GDXr(u0,u1,u2,u3,u4,u5,v0,v1,v2,v3,v4,v5) (u0*v0+u1*v1 + u2*v2+u3*v3 + u4*v4+u5*v5)
13 #define MULT_GDXi(u0,u1,u2,u3,u4,u5,v0,v1,v2,v3,v4,v5) (u0*v1-u1*v0 + u2*v3-u3*v2 + u4*v5-u5*v4)
14 
15 #define EXT_IMG_R(v1r,v1i,v2r,v2i,w1r,w1i,w2r,w2i) (v1r*w2r - v1i*w2i - v2r*w1r + v2i*w1i)
16 #define EXT_IMG_I(v1r,v1i,v2r,v2i,w1r,w1i,w2r,w2i) (- v1r*w2i - v1i*w2r + v2r*w1i + v2i*w1r)
17 
18 
19 //====================================================================
20 void mult_wilson_Meo_dirac(
21  real_t *RESTRICT v2, real_t *RESTRICT u,
22  real_t *RESTRICT v1, real_t *RESTRICT x1,
23  const int ieo, const int jeo,
24  real_t kappa, int *Nsize, int *bc, int iflag)
25 {
26  int Nx = Nsize[0];
27  int Ny = Nsize[1];
28  int Nz = Nsize[2];
29  int Nt = Nsize[3];
30  int Nst = Nx * Ny * Nz * Nt;
31  int Nst_pad = CEIL_NWP(Nst);
32 
33  int nvst = NVC * ND * Nst_pad;
34  int ngst = NDF * Nst_pad * 2 * NDIM;
35 
36  real_t *RESTRICT u_up = u;
37  real_t *RESTRICT u_dn = u;
38 
39 #pragma acc data present(v2[0:nvst], u_up[0:ngst], u_dn[0:ngst], \
40  v1[0:nvst], x1[0:nvst]) \
41  copyin(bc[0:4], kappa, ieo, jeo, iflag, \
42  Nx, Ny, Nz, Nt, Nst, Nst_pad)
43 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
44  {
45  int Nxy = Nx * Ny;
46  int Nxyz = Nx * Ny * Nz;
47 
48 #pragma acc loop gang worker vector
49  for(int site = 0; site < Nst; ++site){
50  int ix = site % Nx;
51  int iy = (site/Nx) % Ny;
52  int iz = (site/(Nxy)) % Nz;
53  int it = site/(Nxyz);
54  int keo = (jeo + iy + iz + it) % 2;
55 
56  real_t bc2;
57 
58  real_t u_0, u_1, u_2, u_3, u_4, u_5;
59  real_t u_6, u_7, u_8, u_9, u10, u11;
60  real_t u12, u13, u14, u15, u16, u17;
61  real_t vt1_0, vt1_1, vt1_2, vt1_3, vt1_4, vt1_5;
62  real_t vt2_0, vt2_1, vt2_2, vt2_3, vt2_4, vt2_5;
63  real_t wt1r, wt1i, wt2r, wt2i;
64 
65  real_t v2_01, v2_11, v2_21, v2_31, v2_41, v2_51;
66  real_t v2_02, v2_12, v2_22, v2_32, v2_42, v2_52;
67  real_t v2_03, v2_13, v2_23, v2_33, v2_43, v2_53;
68  real_t v2_04, v2_14, v2_24, v2_34, v2_44, v2_54;
69 
70 #include "inc/mult_Wilson_eo_xyz_openacc-inc.h"
71 
72 #include "inc/mult_Wilson_eo_t_dirac_openacc-inc.h"
73 
74  // aypx and write back to global memory
75  if(iflag == 0){
76  v2[IDX2_SP(0, 0, site)] = - kappa * v2_01;
77  v2[IDX2_SP(1, 0, site)] = - kappa * v2_11;
78  v2[IDX2_SP(2, 0, site)] = - kappa * v2_21;
79  v2[IDX2_SP(3, 0, site)] = - kappa * v2_31;
80  v2[IDX2_SP(4, 0, site)] = - kappa * v2_41;
81  v2[IDX2_SP(5, 0, site)] = - kappa * v2_51;
82 
83  v2[IDX2_SP(0, 1, site)] = - kappa * v2_02;
84  v2[IDX2_SP(1, 1, site)] = - kappa * v2_12;
85  v2[IDX2_SP(2, 1, site)] = - kappa * v2_22;
86  v2[IDX2_SP(3, 1, site)] = - kappa * v2_32;
87  v2[IDX2_SP(4, 1, site)] = - kappa * v2_42;
88  v2[IDX2_SP(5, 1, site)] = - kappa * v2_52;
89 
90  v2[IDX2_SP(0, 2, site)] = - kappa * v2_03;
91  v2[IDX2_SP(1, 2, site)] = - kappa * v2_13;
92  v2[IDX2_SP(2, 2, site)] = - kappa * v2_23;
93  v2[IDX2_SP(3, 2, site)] = - kappa * v2_33;
94  v2[IDX2_SP(4, 2, site)] = - kappa * v2_43;
95  v2[IDX2_SP(5, 2, site)] = - kappa * v2_53;
96 
97  v2[IDX2_SP(0, 3, site)] = - kappa * v2_04;
98  v2[IDX2_SP(1, 3, site)] = - kappa * v2_14;
99  v2[IDX2_SP(2, 3, site)] = - kappa * v2_24;
100  v2[IDX2_SP(3, 3, site)] = - kappa * v2_34;
101  v2[IDX2_SP(4, 3, site)] = - kappa * v2_44;
102  v2[IDX2_SP(5, 3, site)] = - kappa * v2_54;
103  }else{
104  v2[IDX2_SP(0, 0, site)] = x1[IDX2_SP(0, 0, site)] + kappa * v2_01;
105  v2[IDX2_SP(1, 0, site)] = x1[IDX2_SP(1, 0, site)] + kappa * v2_11;
106  v2[IDX2_SP(2, 0, site)] = x1[IDX2_SP(2, 0, site)] + kappa * v2_21;
107  v2[IDX2_SP(3, 0, site)] = x1[IDX2_SP(3, 0, site)] + kappa * v2_31;
108  v2[IDX2_SP(4, 0, site)] = x1[IDX2_SP(4, 0, site)] + kappa * v2_41;
109  v2[IDX2_SP(5, 0 ,site)] = x1[IDX2_SP(5, 0, site)] + kappa * v2_51;
110 
111  v2[IDX2_SP(0, 1, site)] = x1[IDX2_SP(0, 1, site)] + kappa * v2_02;
112  v2[IDX2_SP(1, 1, site)] = x1[IDX2_SP(1, 1, site)] + kappa * v2_12;
113  v2[IDX2_SP(2, 1, site)] = x1[IDX2_SP(2, 1, site)] + kappa * v2_22;
114  v2[IDX2_SP(3, 1, site)] = x1[IDX2_SP(3, 1, site)] + kappa * v2_32;
115  v2[IDX2_SP(4, 1, site)] = x1[IDX2_SP(4, 1, site)] + kappa * v2_42;
116  v2[IDX2_SP(5, 1, site)] = x1[IDX2_SP(5, 1, site)] + kappa * v2_52;
117 
118  v2[IDX2_SP(0, 2, site)] = x1[IDX2_SP(0, 2, site)] + kappa * v2_03;
119  v2[IDX2_SP(1, 2, site)] = x1[IDX2_SP(1, 2, site)] + kappa * v2_13;
120  v2[IDX2_SP(2, 2, site)] = x1[IDX2_SP(2, 2, site)] + kappa * v2_23;
121  v2[IDX2_SP(3, 2, site)] = x1[IDX2_SP(3, 2, site)] + kappa * v2_33;
122  v2[IDX2_SP(4, 2, site)] = x1[IDX2_SP(4, 2, site)] + kappa * v2_43;
123  v2[IDX2_SP(5, 2, site)] = x1[IDX2_SP(5, 2, site)] + kappa * v2_53;
124 
125  v2[IDX2_SP(0, 3, site)] = x1[IDX2_SP(0, 3, site)] + kappa * v2_04;
126  v2[IDX2_SP(1, 3, site)] = x1[IDX2_SP(1, 3, site)] + kappa * v2_14;
127  v2[IDX2_SP(2, 3, site)] = x1[IDX2_SP(2, 3, site)] + kappa * v2_24;
128  v2[IDX2_SP(3, 3, site)] = x1[IDX2_SP(3, 3, site)] + kappa * v2_34;
129  v2[IDX2_SP(4, 3, site)] = x1[IDX2_SP(4, 3, site)] + kappa * v2_44;
130  v2[IDX2_SP(5, 3, site)] = x1[IDX2_SP(5, 3, site)] + kappa * v2_54;
131  }
132 
133  }
134 
135  }
136 
137 }
138 
139 //====================================================================
140 void mult_wilson_Meo_chiral(
141  real_t *RESTRICT v2, real_t *RESTRICT u,
142  real_t *RESTRICT v1, real_t *RESTRICT x1,
143  const int ieo, const int jeo,
144  real_t kappa, int *Nsize, int *bc, int iflag)
145 {
146  int Nx = Nsize[0];
147  int Ny = Nsize[1];
148  int Nz = Nsize[2];
149  int Nt = Nsize[3];
150  int Nst = Nx * Ny * Nz * Nt;
151  int Nst_pad = CEIL_NWP(Nst);
152 
153  int nvst = NVC * ND * Nst_pad;
154  int ngst = NDF * Nst_pad * 2 * NDIM;
155 
156  real_t *RESTRICT u_up = u;
157  real_t *RESTRICT u_dn = u;
158 
159 #pragma acc data present(v2[0:nvst], u_up[0:ngst], u_dn[0:ngst], \
160  v1[0:nvst], x1[0:nvst]) \
161  copyin(bc[0:4], kappa, ieo, jeo, iflag, \
162  Nx, Ny, Nz, Nt, Nst, Nst_pad)
163 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
164  {
165 
166 #pragma acc loop gang worker vector
167  for(int site = 0; site < Nst; ++site){
168  int ix = site % Nx;
169  int iy = (site/Nx) % Ny;
170  int iz = (site/(Nx*Ny)) % Nz;
171  int it = site/(Nx*Ny*Nz);
172  int keo = (jeo + iy + iz + it) % 2;
173 
174  int Nxy = Nx * Ny;
175  int Nxyz = Nx * Ny * Nz;
176 
177  real_t bc2;
178 
179  real_t u_0, u_1, u_2, u_3, u_4, u_5;
180  real_t u_6, u_7, u_8, u_9, u10, u11;
181  real_t u12, u13, u14, u15, u16, u17;
182  real_t vt1_0, vt1_1, vt1_2, vt1_3, vt1_4, vt1_5;
183  real_t vt2_0, vt2_1, vt2_2, vt2_3, vt2_4, vt2_5;
184  real_t wt1r, wt1i, wt2r, wt2i;
185 
186  real_t v2_01, v2_11, v2_21, v2_31, v2_41, v2_51;
187  real_t v2_02, v2_12, v2_22, v2_32, v2_42, v2_52;
188  real_t v2_03, v2_13, v2_23, v2_33, v2_43, v2_53;
189  real_t v2_04, v2_14, v2_24, v2_34, v2_44, v2_54;
190 
191 #include "inc/mult_Wilson_eo_xyz_openacc-inc.h"
192 
193 #include "inc/mult_Wilson_eo_t_chiral_openacc-inc.h"
194 
195  // aypx and write back to global memory
196  if(iflag == 0){
197  v2[IDX2_SP(0, 0, site)] = - kappa * v2_01;
198  v2[IDX2_SP(1, 0, site)] = - kappa * v2_11;
199  v2[IDX2_SP(2, 0, site)] = - kappa * v2_21;
200  v2[IDX2_SP(3, 0, site)] = - kappa * v2_31;
201  v2[IDX2_SP(4, 0, site)] = - kappa * v2_41;
202  v2[IDX2_SP(5, 0, site)] = - kappa * v2_51;
203 
204  v2[IDX2_SP(0, 1, site)] = - kappa * v2_02;
205  v2[IDX2_SP(1, 1, site)] = - kappa * v2_12;
206  v2[IDX2_SP(2, 1, site)] = - kappa * v2_22;
207  v2[IDX2_SP(3, 1, site)] = - kappa * v2_32;
208  v2[IDX2_SP(4, 1, site)] = - kappa * v2_42;
209  v2[IDX2_SP(5, 1, site)] = - kappa * v2_52;
210 
211  v2[IDX2_SP(0, 2, site)] = - kappa * v2_03;
212  v2[IDX2_SP(1, 2, site)] = - kappa * v2_13;
213  v2[IDX2_SP(2, 2, site)] = - kappa * v2_23;
214  v2[IDX2_SP(3, 2, site)] = - kappa * v2_33;
215  v2[IDX2_SP(4, 2, site)] = - kappa * v2_43;
216  v2[IDX2_SP(5, 2, site)] = - kappa * v2_53;
217 
218  v2[IDX2_SP(0, 3, site)] = - kappa * v2_04;
219  v2[IDX2_SP(1, 3, site)] = - kappa * v2_14;
220  v2[IDX2_SP(2, 3, site)] = - kappa * v2_24;
221  v2[IDX2_SP(3, 3, site)] = - kappa * v2_34;
222  v2[IDX2_SP(4, 3, site)] = - kappa * v2_44;
223  v2[IDX2_SP(5, 3, site)] = - kappa * v2_54;
224  }else{
225  v2[IDX2_SP(0,0,site)] = x1[IDX2_SP(0, 0, site)] + kappa * v2_01;
226  v2[IDX2_SP(1,0,site)] = x1[IDX2_SP(1, 0, site)] + kappa * v2_11;
227  v2[IDX2_SP(2,0,site)] = x1[IDX2_SP(2, 0, site)] + kappa * v2_21;
228  v2[IDX2_SP(3,0,site)] = x1[IDX2_SP(3, 0, site)] + kappa * v2_31;
229  v2[IDX2_SP(4,0,site)] = x1[IDX2_SP(4, 0, site)] + kappa * v2_41;
230  v2[IDX2_SP(5,0,site)] = x1[IDX2_SP(5, 0, site)] + kappa * v2_51;
231 
232  v2[IDX2_SP(0,1,site)] = x1[IDX2_SP(0, 1, site)] + kappa * v2_02;
233  v2[IDX2_SP(1,1,site)] = x1[IDX2_SP(1, 1, site)] + kappa * v2_12;
234  v2[IDX2_SP(2,1,site)] = x1[IDX2_SP(2, 1, site)] + kappa * v2_22;
235  v2[IDX2_SP(3,1,site)] = x1[IDX2_SP(3, 1, site)] + kappa * v2_32;
236  v2[IDX2_SP(4,1,site)] = x1[IDX2_SP(4, 1, site)] + kappa * v2_42;
237  v2[IDX2_SP(5,1,site)] = x1[IDX2_SP(5, 1, site)] + kappa * v2_52;
238 
239  v2[IDX2_SP(0,2,site)] = x1[IDX2_SP(0, 2, site)] + kappa * v2_03;
240  v2[IDX2_SP(1,2,site)] = x1[IDX2_SP(1, 2, site)] + kappa * v2_13;
241  v2[IDX2_SP(2,2,site)] = x1[IDX2_SP(2, 2, site)] + kappa * v2_23;
242  v2[IDX2_SP(3,2,site)] = x1[IDX2_SP(3, 2, site)] + kappa * v2_33;
243  v2[IDX2_SP(4,2,site)] = x1[IDX2_SP(4, 2, site)] + kappa * v2_43;
244  v2[IDX2_SP(5,2,site)] = x1[IDX2_SP(5, 2, site)] + kappa * v2_53;
245 
246  v2[IDX2_SP(0,3,site)] = x1[IDX2_SP(0, 3, site)] + kappa * v2_04;
247  v2[IDX2_SP(1,3,site)] = x1[IDX2_SP(1, 3, site)] + kappa * v2_14;
248  v2[IDX2_SP(2,3,site)] = x1[IDX2_SP(2, 3, site)] + kappa * v2_24;
249  v2[IDX2_SP(3,3,site)] = x1[IDX2_SP(3, 3, site)] + kappa * v2_34;
250  v2[IDX2_SP(4,3,site)] = x1[IDX2_SP(4, 3, site)] + kappa * v2_44;
251  v2[IDX2_SP(5,3,site)] = x1[IDX2_SP(5, 3, site)] + kappa * v2_54;
252  }
253 
254  }
255 
256  }
257 
258 }
259 
260 //====================================================================
261 void mult_wilson_eo_1_dirac(
262  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
263  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
264  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
265  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
266  real_t *RESTRICT u, real_t *RESTRICT v1,
267  const int ieo, const int jeo,
268  int *Nsize, int *bc, int *do_comm, int Nc)
269 {
270  int Nx = Nsize[0];
271  int Ny = Nsize[1];
272  int Nz = Nsize[2];
273  int Nt = Nsize[3];
274  int Nst = Nx * Ny * Nz * Nt;
275  int Nst_pad = CEIL_NWP(Nst);
276 
277  int size = NVC * ND * Nst_pad;
278  int size_u = NDF * Nst_pad * 2 * NDIM;
279  int size_bx = NVC * ND2 * CEIL_NWP((Ny * Nz * Nt + 1)/2);
280  int size_by = NVC * ND2 * CEIL_NWP(Nx * Nz * Nt);
281  int size_bz = NVC * ND2 * CEIL_NWP(Nx * Ny * Nt);
282  int size_bt = NVC * ND2 * CEIL_NWP(Nx * Ny * Nz);
283 
284 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
285  buf_yp[0:size_by], buf_ym[0:size_by], \
286  buf_zp[0:size_bz], buf_zm[0:size_bz], \
287  buf_tp[0:size_bt], buf_tm[0:size_bt], \
288  u[0:size_u], v1[0:size]) \
289  copyin(bc[0:4], do_comm[0:4], ieo, jeo, \
290  Nst, Nst_pad, Nx, Ny, Nz, Nt)
291 {
292 
293  // int idir = 0;
294  if(do_comm[0] > 0){
295 
296 #pragma acc parallel async \
297  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
298  {
299  int Nyzt = Ny * Nz * Nt;
300 
301 #pragma acc loop gang worker vector
302  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
303  int ix = 0;
304  int iy = iyzt % Ny;
305  int iz = (iyzt/Ny) % Nz;
306  int it = iyzt/(Ny * Nz);
307  int ist = ix + Nx * iyzt;
308  int keo = (jeo + iy + iz + it) % 2;
309 
310  if(keo == 1){
311  int iyzt2 = iyzt/2;
312  real_t bc2 = bc[0];
313  for(int ic = 0; ic < NC; ++ic){
314  real_t vt1[2], vt2[2];
315  vt1[0] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
316  vt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
317  vt2[0] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
318  vt2[1] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
319  buf_xp[IDXBF_R(ic, 0, iyzt2)] = bc2 * vt1[0];
320  buf_xp[IDXBF_I(ic, 0, iyzt2)] = bc2 * vt1[1];
321  buf_xp[IDXBF_R(ic, 1, iyzt2)] = bc2 * vt2[0];
322  buf_xp[IDXBF_I(ic, 1, iyzt2)] = bc2 * vt2[1];
323  }
324  }
325 
326  }
327  }
328 
329 #pragma acc parallel async \
330  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
331  {
332  int Nyzt = Ny * Nz * Nt;
333 
334 #pragma acc loop gang worker vector
335  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
336  int ix = Nx-1;
337  int iy = iyzt % Ny;
338  int iz = (iyzt/Ny) % Nz;
339  int it = iyzt/(Ny * Nz);
340  int ist = ix + Nx * iyzt;
341  int istu = ist + Nst_pad * (1-ieo + 2*0);
342  int keo = (jeo + iy + iz + it) % 2;
343 
344  if(keo == 0){
345  int iyzt2 = iyzt/2;
346 
347  real_t vt1[NVC], vt2[NVC], ut[NDF];
348  for(int ic = 0; ic < NC; ++ic){
349  int icr = 2*ic;
350  int ici = 2*ic + 1;
351  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
352  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
353  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
354  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
355  }
356 
357  real_t wt1[2], wt2[2];
358 
359  for(int ic = 0; ic < NC; ++ic){
360 
361  for(int ic2 = 0; ic2 < NC; ++ic2){
362  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
363  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
364  }
365 
366  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
367  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
368  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
369  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
370  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
371  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
372  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
373  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
374 
375  buf_xm[IDXBF_R(ic, 0, iyzt2)] = wt1[0];
376  buf_xm[IDXBF_I(ic, 0, iyzt2)] = wt1[1];
377  buf_xm[IDXBF_R(ic, 1, iyzt2)] = wt2[0];
378  buf_xm[IDXBF_I(ic, 1, iyzt2)] = wt2[1];
379  }
380 
381  }
382 
383  }
384 
385  }
386 
387  } // do_comm[0]
388 
389  // idir = 1;
390  if(do_comm[1] > 0){
391 
392 #pragma acc parallel async \
393  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
394  {
395  int Nzt = Nz * Nt;
396 
397 #pragma acc loop gang worker vector collapse(2)
398  for(int izt = 0; izt < Nzt; ++izt){
399  for(int ix = 0; ix < Nx; ++ix){
400  int iy = 0;
401  int ist = ix + Nx * (iy + Ny * izt);
402  int ixzt = ix + Nx * izt;
403  real_t bc2 = bc[1];
404 
405  for(int ic = 0; ic < NC; ++ic){
406  real_t vt1[2], vt2[2];
407  vt1[0] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
408  vt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
409  vt2[0] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
410  vt2[1] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
411  buf_yp[IDXBF_R(ic, 0, ixzt)] = bc2 * vt1[0];
412  buf_yp[IDXBF_I(ic, 0, ixzt)] = bc2 * vt1[1];
413  buf_yp[IDXBF_R(ic, 1, ixzt)] = bc2 * vt2[0];
414  buf_yp[IDXBF_I(ic, 1, ixzt)] = bc2 * vt2[1];
415  }
416 
417  }
418  }
419 
420  }
421 
422 #pragma acc parallel async \
423  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
424  {
425  int Nzt = Nz * Nt;
426 
427 #pragma acc loop gang worker vector collapse(2)
428  for(int izt = 0; izt < Nzt; ++izt){
429  for(int ix = 0; ix < Nx; ++ix){
430  int iy = Ny-1;
431  int ist = ix + Nx * (iy + Ny*izt);
432  int istu = ist + Nst_pad * (1-ieo + 2*1);
433  int ixzt = ix + Nx * izt;
434 
435  real_t vt1[NVC], vt2[NVC];
436  for(int ic = 0; ic < NC; ++ic){
437  int icr = 2*ic;
438  int ici = 2*ic + 1;
439  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
440  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
441  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
442  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
443  }
444 
445  for(int ic = 0; ic < NC; ++ic){
446 
447  real_t ut[NVC];
448  for(int ic2 = 0; ic2 < NC; ++ic2){
449  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
450  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
451  }
452 
453  real_t wt1[2], wt2[2];
454  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
455  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
456  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
457  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
458  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
459  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
460  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
461  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
462 
463  buf_ym[IDXBF_R(ic, 0, ixzt)] = wt1[0];
464  buf_ym[IDXBF_I(ic, 0, ixzt)] = wt1[1];
465  buf_ym[IDXBF_R(ic, 1, ixzt)] = wt2[0];
466  buf_ym[IDXBF_I(ic, 1, ixzt)] = wt2[1];
467 
468  }
469 
470  }
471  }
472 
473  }
474 
475  } // do_comm[1]
476 
477  // idir = 2;
478  if(do_comm[2] > 0){
479 
480 #pragma acc parallel async \
481  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
482  {
483  int Nxy = Nx * Ny;
484 
485 #pragma acc loop gang worker vector collapse(2)
486  for(int it = 0; it < Nt; ++it){
487  for(int ixy = 0; ixy < Nxy; ++ixy){
488  int iz = 0;
489  int ist = ixy + Nxy * (iz + Nz * it);
490  int ixyt = ixy + Nxy * it;
491  real_t bc2 = bc[2];
492 
493  for(int ic = 0; ic < NC; ++ic){
494  real_t wt1[2], wt2[2];
495  wt1[0] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
496  wt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
497  wt2[0] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
498  wt2[1] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
499  buf_zp[IDXBF_R(ic, 0, ixyt)] = bc2 * wt1[0];
500  buf_zp[IDXBF_I(ic, 0, ixyt)] = bc2 * wt1[1];
501  buf_zp[IDXBF_R(ic, 1, ixyt)] = bc2 * wt2[0];
502  buf_zp[IDXBF_I(ic, 1, ixyt)] = bc2 * wt2[1];
503  }
504 
505  }
506  }
507 
508  }
509 
510 #pragma acc parallel async \
511  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
512  {
513  int Nxy = Nx * Ny;
514 
515 #pragma acc loop gang worker vector collapse(2)
516  for(int it = 0; it < Nt; ++it){
517  for(int ixy = 0; ixy < Nxy; ++ixy){
518  int iz = Nz-1;
519  int ist = ixy + Nxy * (iz + Nz * it);
520  int istu = ist + Nst_pad * (1-ieo + 2*2);
521  int ixyt = ixy + Nxy * it;
522 
523  real_t vt1[NVC], vt2[NVC];
524  for(int ic = 0; ic < NC; ++ic){
525  int icr = 2*ic;
526  int ici = 2*ic + 1;
527  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
528  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
529  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
530  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
531  }
532 
533  for(int ic = 0; ic < NC; ++ic){
534 
535  real_t ut[NVC];
536  for(int ic2 = 0; ic2 < NC; ++ic2){
537  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
538  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
539  }
540 
541  real_t wt1[2], wt2[2];
542  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
543  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
544  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
545  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
546  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
547  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
548  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
549  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
550  buf_zm[IDXBF_R(ic, 0, ixyt)] = wt1[0];
551  buf_zm[IDXBF_I(ic, 0, ixyt)] = wt1[1];
552  buf_zm[IDXBF_R(ic, 1, ixyt)] = wt2[0];
553  buf_zm[IDXBF_I(ic, 1, ixyt)] = wt2[1];
554 
555  }
556 
557  }
558  }
559 
560  }
561 
562  } // do_comm[2]
563 
564  // idir = 3;
565  if(do_comm[3] > 0){
566 
567 #pragma acc parallel async \
568  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
569  {
570  int Nxyz = Nx * Ny * Nz;
571 
572 #pragma acc loop gang worker vector
573  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
574  int it = 0;
575  int ist = ixyz + Nxyz * it;
576  real_t bc2 = 2.0 * bc[3];
577  for(int ic = 0; ic < NC; ++ic){
578  buf_tp[IDXBF_R(ic, 0, ixyz)] = bc2 * v1[IDX2_SP_R(ic, 2, ist)];
579  buf_tp[IDXBF_I(ic, 0, ixyz)] = bc2 * v1[IDX2_SP_I(ic, 2, ist)];
580  buf_tp[IDXBF_R(ic, 1, ixyz)] = bc2 * v1[IDX2_SP_R(ic, 3, ist)];
581  buf_tp[IDXBF_I(ic, 1, ixyz)] = bc2 * v1[IDX2_SP_I(ic, 3, ist)];
582  }
583  }
584 
585  }
586 
587 #pragma acc parallel async \
588  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
589  {
590  int Nxyz = Nx * Ny * Nz;
591 
592 #pragma acc loop gang worker vector
593  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
594  int it = Nt-1;
595  int ist = ixyz + Nxyz * it;
596  int istu = ist + Nst_pad * (1-ieo + 2*3);
597 
598  real_t vt1[NVC], vt2[NVC];
599  for(int ic = 0; ic < NC; ++ic){
600  int icr = 2*ic;
601  int ici = 2*ic + 1;
602  vt1[icr] = 2.0 * v1[IDX2_SP_R(ic, 0, ist)];
603  vt1[ici] = 2.0 * v1[IDX2_SP_I(ic, 0, ist)];
604  vt2[icr] = 2.0 * v1[IDX2_SP_R(ic, 1, ist)];
605  vt2[ici] = 2.0 * v1[IDX2_SP_I(ic, 1, ist)];
606  }
607 
608  for(int ic = 0; ic < NC; ++ic){
609 
610  real_t ut[NVC];
611  for(int ic2 = 0; ic2 < NC; ++ic2){
612  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
613  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
614  }
615 
616  real_t wt1[2], wt2[2];
617  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
618  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
619  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
620  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
621  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
622  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
623  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
624  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
625  buf_tm[IDXBF_R(ic, 0, ixyz)] = wt1[0];
626  buf_tm[IDXBF_I(ic, 0, ixyz)] = wt1[1];
627  buf_tm[IDXBF_R(ic, 1, ixyz)] = wt2[0];
628  buf_tm[IDXBF_I(ic, 1, ixyz)] = wt2[1];
629  }
630 
631  }
632 
633  }
634 
635  } // do_comm[3]
636 
637  #pragma acc wait
638 
639  } // acc data
640 
641  if(do_comm[0] > 0){
642  #pragma acc update async host (buf_xp[0:size_bx])
643  #pragma acc update async host (buf_xm[0:size_bx])
644  }
645  if(do_comm[1] > 0){
646  #pragma acc update async host (buf_yp[0:size_by])
647  #pragma acc update async host (buf_ym[0:size_by])
648  }
649  if(do_comm[2] > 0){
650  #pragma acc update async host (buf_zp[0:size_bz])
651  #pragma acc update async host (buf_zm[0:size_bz])
652  }
653  if(do_comm[3] > 0){
654  #pragma acc update async host (buf_tp[0:size_bt])
655  #pragma acc update async host (buf_tm[0:size_bt])
656  }
657 
658  #pragma acc wait
659 
660 }
661 
662 //====================================================================
663 void mult_wilson_eo_1_chiral(
664  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
665  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
666  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
667  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
668  real_t *RESTRICT u, real_t *RESTRICT v1,
669  const int ieo, const int jeo,
670  int *Nsize, int *bc, int *do_comm, int Nc)
671 {
672  int Nx = Nsize[0];
673  int Ny = Nsize[1];
674  int Nz = Nsize[2];
675  int Nt = Nsize[3];
676  int Nst = Nx * Ny * Nz * Nt;
677  int Nst_pad = CEIL_NWP(Nst);
678 
679  int size = NVC * ND * Nst_pad;
680  int size_u = NDF * 2 * Nst_pad * NDIM;
681  int size_bx = NVC * ND2 * CEIL_NWP((Ny * Nz * Nt + 1)/2);
682  int size_by = NVC * ND2 * CEIL_NWP(Nx * Nz * Nt);
683  int size_bz = NVC * ND2 * CEIL_NWP(Nx * Ny * Nt);
684  int size_bt = NVC * ND2 * CEIL_NWP(Nx * Ny * Nz);
685 
686 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
687  buf_yp[0:size_by], buf_ym[0:size_by], \
688  buf_zp[0:size_bz], buf_zm[0:size_bz], \
689  buf_tp[0:size_bt], buf_tm[0:size_bt], \
690  u[0:size_u], v1[0:size]) \
691  copyin(bc[0:4], do_comm[0:4], ieo, jeo, \
692  Nx, Ny, Nz, Nt, Nst, Nst_pad)
693 {
694 
695  // int idir = 0;
696  if(do_comm[0] > 0){
697 
698 #pragma acc parallel async \
699  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
700  {
701  int Nyzt = Ny * Nz * Nt;
702 
703 #pragma acc loop gang worker vector
704  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
705  int ix = 0;
706  int iy = iyzt % Ny;
707  int iz = (iyzt/Ny) % Nz;
708  int it = iyzt/(Ny * Nz);
709  int ist = ix + Nx * iyzt;
710  int keo = (jeo + iy + iz + it) % 2;
711 
712  if(keo == 1){
713  int iyzt2 = iyzt/2;
714  real_t bc2 = bc[0];
715  for(int ic = 0; ic < NC; ++ic){
716  real_t vt1[2], vt2[2];
717  vt1[0] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
718  vt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
719  vt2[0] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
720  vt2[1] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
721  buf_xp[IDXBF_R(ic, 0, iyzt2)] = bc2 * vt1[0];
722  buf_xp[IDXBF_I(ic, 0, iyzt2)] = bc2 * vt1[1];
723  buf_xp[IDXBF_R(ic, 1, iyzt2)] = bc2 * vt2[0];
724  buf_xp[IDXBF_I(ic, 1, iyzt2)] = bc2 * vt2[1];
725  }
726  }
727 
728  }
729  }
730 
731 #pragma acc parallel async \
732  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
733  {
734  int Nyzt = Ny * Nz * Nt;
735 
736 #pragma acc loop gang worker vector
737  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
738  int ix = Nx-1;
739  int iy = iyzt % Ny;
740  int iz = (iyzt/Ny) % Nz;
741  int it = iyzt/(Ny * Nz);
742  int ist = ix + Nx * iyzt;
743  int istu = ist + Nst_pad * (1-ieo + 2*0);
744  int keo = (jeo + iy + iz + it) % 2;
745 
746  if(keo == 0){
747  int iyzt2 = iyzt/2;
748 
749  real_t vt1[NVC], vt2[NVC], ut[NDF];
750  for(int ic = 0; ic < NC; ++ic){
751  int icr = 2*ic;
752  int ici = 2*ic + 1;
753  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
754  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
755  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
756  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
757  }
758 
759  real_t wt1[2], wt2[2];
760 
761  for(int ic = 0; ic < NC; ++ic){
762 
763  for(int ic2 = 0; ic2 < NC; ++ic2){
764  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
765  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
766  }
767 
768  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
769  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
770  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
771  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
772  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
773  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
774  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
775  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
776 
777  buf_xm[IDXBF_R(ic, 0, iyzt2)] = wt1[0];
778  buf_xm[IDXBF_I(ic, 0, iyzt2)] = wt1[1];
779  buf_xm[IDXBF_R(ic, 1, iyzt2)] = wt2[0];
780  buf_xm[IDXBF_I(ic, 1, iyzt2)] = wt2[1];
781  }
782 
783  }
784 
785  }
786 
787  }
788 
789  } // do_comm[0]
790 
791  // idir = 1;
792  if(do_comm[1] > 0){
793 
794 #pragma acc parallel async \
795  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
796  {
797  int Nzt = Nz * Nt;
798 
799 #pragma acc loop gang worker vector collapse(2)
800  for(int izt = 0; izt < Nzt; ++izt){
801  for(int ix = 0; ix < Nx; ++ix){
802  int iy = 0;
803  int ist = ix + Nx * (iy + Ny * izt);
804  int ixzt = ix + Nx * izt;
805  real_t bc2 = bc[1];
806 
807  for(int ic = 0; ic < NC; ++ic){
808  real_t vt1[2], vt2[2];
809  vt1[0] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
810  vt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
811  vt2[0] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
812  vt2[1] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
813  buf_yp[IDXBF_R(ic, 0, ixzt)] = bc2 * vt1[0];
814  buf_yp[IDXBF_I(ic, 0, ixzt)] = bc2 * vt1[1];
815  buf_yp[IDXBF_R(ic, 1, ixzt)] = bc2 * vt2[0];
816  buf_yp[IDXBF_I(ic, 1, ixzt)] = bc2 * vt2[1];
817  }
818 
819  }
820  }
821 
822  }
823 
824 #pragma acc parallel async \
825  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
826  {
827  int Nzt = Nz * Nt;
828 
829 #pragma acc loop gang worker vector collapse(2)
830  for(int izt = 0; izt < Nzt; ++izt){
831  for(int ix = 0; ix < Nx; ++ix){
832  int iy = Ny-1;
833  int ist = ix + Nx * (iy + Ny*izt);
834  int istu = ist + Nst_pad * (1-ieo + 2*1);
835  int ixzt = ix + Nx * izt;
836 
837  real_t vt1[NVC], vt2[NVC];
838  for(int ic = 0; ic < NC; ++ic){
839  int icr = 2*ic;
840  int ici = 2*ic + 1;
841  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
842  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
843  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
844  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
845  }
846 
847  for(int ic = 0; ic < NC; ++ic){
848 
849  real_t ut[NVC];
850  for(int ic2 = 0; ic2 < NC; ++ic2){
851  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
852  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
853  }
854 
855  real_t wt1[2], wt2[2];
856  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
857  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
858  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
859  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
860  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
861  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
862  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
863  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
864 
865  buf_ym[IDXBF_R(ic, 0, ixzt)] = wt1[0];
866  buf_ym[IDXBF_I(ic, 0, ixzt)] = wt1[1];
867  buf_ym[IDXBF_R(ic, 1, ixzt)] = wt2[0];
868  buf_ym[IDXBF_I(ic, 1, ixzt)] = wt2[1];
869 
870  }
871 
872  }
873  }
874 
875  }
876 
877  } // do_comm[1]
878 
879  // idir = 2;
880  if(do_comm[2] > 0){
881 
882 #pragma acc parallel async \
883  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
884  {
885  int Nxy = Nx * Ny;
886 
887 #pragma acc loop gang worker vector collapse(2)
888  for(int it = 0; it < Nt; ++it){
889  for(int ixy = 0; ixy < Nxy; ++ixy){
890  int iz = 0;
891  int ist = ixy + Nxy * (iz + Nz * it);
892  int ixyt = ixy + Nxy * it;
893  real_t bc2 = bc[2];
894 
895  for(int ic = 0; ic < NC; ++ic){
896  real_t wt1[2], wt2[2];
897  wt1[0] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
898  wt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
899  wt2[0] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
900  wt2[1] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
901  buf_zp[IDXBF_R(ic, 0, ixyt)] = bc2 * wt1[0];
902  buf_zp[IDXBF_I(ic, 0, ixyt)] = bc2 * wt1[1];
903  buf_zp[IDXBF_R(ic, 1, ixyt)] = bc2 * wt2[0];
904  buf_zp[IDXBF_I(ic, 1, ixyt)] = bc2 * wt2[1];
905  }
906 
907  }
908  }
909 
910  }
911 
912 #pragma acc parallel async \
913  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
914  {
915  int Nxy = Nx * Ny;
916 
917 #pragma acc loop gang worker vector collapse(2)
918  for(int it = 0; it < Nt; ++it){
919  for(int ixy = 0; ixy < Nxy; ++ixy){
920  int iz = Nz-1;
921  int ist = ixy + Nxy * (iz + Nz * it);
922  int istu = ist + Nst_pad * (1-ieo + 2*2);
923  int ixyt = ixy + Nxy * it;
924 
925  real_t vt1[NVC], vt2[NVC];
926  for(int ic = 0; ic < NC; ++ic){
927  int icr = 2*ic;
928  int ici = 2*ic + 1;
929  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
930  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
931  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
932  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
933  }
934 
935  for(int ic = 0; ic < NC; ++ic){
936 
937  real_t ut[NVC];
938  for(int ic2 = 0; ic2 < NC; ++ic2){
939  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
940  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
941  }
942 
943  real_t wt1[2], wt2[2];
944  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
945  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
946  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
947  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
948  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
949  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
950  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
951  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
952  buf_zm[IDXBF_R(ic, 0, ixyt)] = wt1[0];
953  buf_zm[IDXBF_I(ic, 0, ixyt)] = wt1[1];
954  buf_zm[IDXBF_R(ic, 1, ixyt)] = wt2[0];
955  buf_zm[IDXBF_I(ic, 1, ixyt)] = wt2[1];
956 
957  }
958 
959  }
960  }
961 
962  }
963 
964  } // do_comm[2]
965 
966  // idir = 3;
967  if(do_comm[3] > 0){
968 
969 #pragma acc parallel async\
970  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
971  {
972  int Nxyz = Nx * Ny * Nz;
973 
974 #pragma acc loop gang worker vector
975  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
976  int it = 0;
977  int ist = ixyz + Nxyz * it;
978  real_t bc2 = bc[3];
979  for(int ivc = 0; ivc < NVC; ++ivc){
980  buf_tp[IDXBF(ivc, 0, ixyz)]
981  = bc2 * (v1[IDX2_SP(ivc, 0, ist)] + v1[IDX2_SP(ivc, 2, ist)]);
982  buf_tp[IDXBF(ivc, 1, ixyz)]
983  = bc2 * (v1[IDX2_SP(ivc, 1, ist)] + v1[IDX2_SP(ivc, 3, ist)]);
984  }
985  }
986 
987  }
988 
989 #pragma acc parallel async \
990  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
991  {
992  int Nxyz = Nx * Ny * Nz;
993 
994 #pragma acc loop gang worker vector
995  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
996  int it = Nt-1;
997  int ist = ixyz + Nxyz * it;
998  int istu = ist + Nst_pad * (1-ieo + 2*3);
999 
1000  real_t vt1[NVC], vt2[NVC];
1001  for(int ivc = 0; ivc < NVC; ++ivc){
1002  vt1[ivc] = v1[IDX2_SP(ivc, 0, ist)] - v1[IDX2_SP(ivc, 2, ist)];
1003  vt2[ivc] = v1[IDX2_SP(ivc, 1, ist)] - v1[IDX2_SP(ivc, 3, ist)];
1004  }
1005 
1006  for(int ic = 0; ic < NC; ++ic){
1007 
1008  real_t ut[NVC];
1009  for(int ic2 = 0; ic2 < NC; ++ic2){
1010  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
1011  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
1012  }
1013 
1014  real_t wt1[2], wt2[2];
1015  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1016  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1017  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1018  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1019  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1020  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1021  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1022  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1023  buf_tm[IDXBF_R(ic, 0, ixyz)] = wt1[0];
1024  buf_tm[IDXBF_I(ic, 0, ixyz)] = wt1[1];
1025  buf_tm[IDXBF_R(ic, 1, ixyz)] = wt2[0];
1026  buf_tm[IDXBF_I(ic, 1, ixyz)] = wt2[1];
1027 
1028  }
1029 
1030  }
1031 
1032  }
1033 
1034  } // do_comm[3]
1035 
1036  #pragma acc wait
1037 
1038  } // acc data
1039 
1040  if(do_comm[0] > 0){
1041  #pragma acc update async host (buf_xp[0:size_bx])
1042  #pragma acc update async host (buf_xm[0:size_bx])
1043  }
1044  if(do_comm[1] > 0){
1045  #pragma acc update async host (buf_yp[0:size_by])
1046  #pragma acc update async host (buf_ym[0:size_by])
1047  }
1048  if(do_comm[2] > 0){
1049  #pragma acc update async host (buf_zp[0:size_bz])
1050  #pragma acc update async host (buf_zm[0:size_bz])
1051  }
1052  if(do_comm[3] > 0){
1053  #pragma acc update async host (buf_tp[0:size_bt])
1054  #pragma acc update async host (buf_tm[0:size_bt])
1055  }
1056 
1057  #pragma acc wait
1058 
1059 }
1060 
1061 //====================================================================
1062 void mult_wilson_eo_2_dirac(real_t *RESTRICT v2, real_t *RESTRICT u,
1063  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
1064  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
1065  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
1066  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
1067  real_t kappa, const int ieo, const int jeo,
1068  const int iflag,
1069  int *Nsize, int *bc, int *do_comm, int Nc)
1070 {
1071  int Nx = Nsize[0];
1072  int Ny = Nsize[1];
1073  int Nz = Nsize[2];
1074  int Nt = Nsize[3];
1075  int Nst = Nx * Ny * Nz * Nt;
1076  int Nst_pad = CEIL_NWP(Nst);
1077 
1078  int size = NVC * ND * Nst_pad;
1079  int size_u = NDF * Nst_pad * 2 * NDIM;
1080  int size_bx = NVC * ND2 * CEIL_NWP((Ny * Nz * Nt + 1)/2);
1081  int size_by = NVC * ND2 * CEIL_NWP(Nx * Nz * Nt);
1082  int size_bz = NVC * ND2 * CEIL_NWP(Nx * Ny * Nt);
1083  int size_bt = NVC * ND2 * CEIL_NWP(Nx * Ny * Nz);
1084 
1085  if(do_comm[0] > 0){
1086  #pragma acc update async device (buf_xp[0:size_bx])
1087  #pragma acc update async device (buf_xm[0:size_bx])
1088  }
1089  if(do_comm[1] > 0){
1090  #pragma acc update async device (buf_yp[0:size_by])
1091  #pragma acc update async device (buf_ym[0:size_by])
1092  }
1093  if(do_comm[2] > 0){
1094  #pragma acc update async device (buf_zp[0:size_bz])
1095  #pragma acc update async device (buf_zm[0:size_bz])
1096  }
1097  if(do_comm[3] > 0){
1098  #pragma acc update async device (buf_tp[0:size_bt])
1099  #pragma acc update async device (buf_tm[0:size_bt])
1100  }
1101 
1102  #pragma acc wait
1103 
1104 
1105 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
1106  buf_yp[0:size_by], buf_ym[0:size_by], \
1107  buf_zp[0:size_bz], buf_zm[0:size_bz], \
1108  buf_tp[0:size_bt], buf_tm[0:size_bt], \
1109  v2[0:size], u[0:size_u]) \
1110  copyin(bc[0:4], do_comm[0:4], ieo, jeo, kappa, \
1111  Nx, Ny, Nz, Nt, Nst, Nst_pad)
1112 {
1113 
1114 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1115  {
1116 #pragma acc loop gang worker vector
1117  for(int ist = 0; ist < Nst; ++ist){
1118  int ix = ist % Nx;
1119  int iy = (ist/Nx) % Ny;
1120  int iz = (ist/(Nx*Ny)) % Nz;
1121  int it = ist/(Nx*Ny*Nz);
1122  int keo = (jeo + iy + iz + it) % 2;
1123 
1124  real_t v2L[NVC*ND];
1125  for(int id = 0; id < ND; ++id){
1126  for(int ivc = 0; ivc < NVC; ++ivc){
1127  v2L[ivc + NVC * id] = 0.0;
1128  }
1129  }
1130  int opr_any = 0;
1131 
1132 #include "inc/mult_Wilson_eo_2xyz_openacc-inc.h"
1133 
1134  // idir = 3
1135  if(do_comm[3] > 0){
1136 
1137  if(it == Nt-1){
1138  int ixyz = ix + Nx * (iy + Ny * iz);
1139  int istu = ist + Nst_pad * (ieo + 2*3);
1140 
1141  real_t vt1[NVC], vt2[NVC];
1142 
1143  for(int ivc = 0; ivc < NVC; ++ivc){
1144  vt1[ivc] = buf_tp[IDXBF(ivc, 0, ixyz)];
1145  vt2[ivc] = buf_tp[IDXBF(ivc, 1, ixyz)];
1146  }
1147 
1148  for(int ic = 0; ic < NC; ++ic){
1149  real_t ut[NVC];
1150  for(int ic2 = 0; ic2 < NC; ++ic2){
1151  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, istu)];
1152  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, istu)];
1153  }
1154  real_t wt1[2], wt2[2];
1155  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1156  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1157  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1158  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1159  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1160  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1161  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1162  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1163  v2L[ic*2 + ID3] += wt1[0];
1164  v2L[ic*2+1 + ID3] += wt1[1];
1165  v2L[ic*2 + ID4] += wt2[0];
1166  v2L[ic*2+1 + ID4] += wt2[1];
1167  }
1168  ++opr_any;
1169  }
1170 
1171  if(it == 0){
1172  int ixyz = ix + Nx * (iy + Ny * iz);
1173  real_t bc2 = bc[3];
1174  real_t wt1[2], wt2[2];
1175 
1176  for(int ic = 0; ic < NC; ++ic){
1177  wt1[0] = bc2 * buf_tm[IDXBF_R(ic, 0, ixyz)];
1178  wt1[1] = bc2 * buf_tm[IDXBF_I(ic, 0, ixyz)];
1179  wt2[0] = bc2 * buf_tm[IDXBF_R(ic, 1, ixyz)];
1180  wt2[1] = bc2 * buf_tm[IDXBF_I(ic, 1, ixyz)];
1181  v2L[ic*2 + ID1] += wt1[0];
1182  v2L[ic*2+1 + ID1] += wt1[1];
1183  v2L[ic*2 + ID2] += wt2[0];
1184  v2L[ic*2+1 + ID2] += wt2[1];
1185  }
1186  ++opr_any;
1187  }
1188 
1189  }
1190 
1191  //axpy
1192  if(opr_any > 0){
1193  if(iflag == 0){
1194  for(int id = 0; id < ND; ++id){
1195  for(int ivc = 0; ivc < NVC; ++ivc){
1196  v2[IDX2_SP(ivc, id, ist)] += -kappa * v2L[ivc + NVC * id];
1197  }
1198  }
1199  }else{
1200  for(int id = 0; id < ND; ++id){
1201  for(int ivc = 0; ivc < NVC; ++ivc){
1202  v2[IDX2_SP(ivc, id, ist)] += kappa * v2L[ivc + NVC * id];
1203  }
1204  }
1205  }
1206  }
1207  }
1208 
1209  }
1210 
1211  }
1212 
1213 }
1214 
1215 //====================================================================
1216 void mult_wilson_eo_2_chiral(real_t *RESTRICT v2, real_t *RESTRICT u,
1217  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
1218  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
1219  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
1220  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
1221  real_t kappa, const int ieo, const int jeo,
1222  const int iflag,
1223  int *Nsize, int *bc, int *do_comm, int Nc)
1224 {
1225  int Nx = Nsize[0];
1226  int Ny = Nsize[1];
1227  int Nz = Nsize[2];
1228  int Nt = Nsize[3];
1229  int Nst = Nx * Ny * Nz * Nt;
1230  int Nst_pad = CEIL_NWP(Nst);
1231 
1232  int size = NVC * ND * Nst_pad;
1233  int size_u = NDF * Nst_pad * 2 * NDIM;
1234  int size_bx = NVC * ND2 * CEIL_NWP((Ny * Nz * Nt + 1)/2);
1235  int size_by = NVC * ND2 * CEIL_NWP(Nx * Nz * Nt);
1236  int size_bz = NVC * ND2 * CEIL_NWP(Nx * Ny * Nt);
1237  int size_bt = NVC * ND2 * CEIL_NWP(Nx * Ny * Nz);
1238 
1239  if(do_comm[0] > 0){
1240  #pragma acc update async device (buf_xp[0:size_bx])
1241  #pragma acc update async device (buf_xm[0:size_bx])
1242  }
1243  if(do_comm[1] > 0){
1244  #pragma acc update async device (buf_yp[0:size_by])
1245  #pragma acc update async device (buf_ym[0:size_by])
1246  }
1247  if(do_comm[2] > 0){
1248  #pragma acc update async device (buf_zp[0:size_bz])
1249  #pragma acc update async device (buf_zm[0:size_bz])
1250  }
1251  if(do_comm[3] > 0){
1252  #pragma acc update async device (buf_tp[0:size_bt])
1253  #pragma acc update async device (buf_tm[0:size_bt])
1254  }
1255 
1256  #pragma acc wait
1257 
1258 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
1259  buf_yp[0:size_by], buf_ym[0:size_by], \
1260  buf_zp[0:size_bz], buf_zm[0:size_bz], \
1261  buf_tp[0:size_bt], buf_tm[0:size_bt], \
1262  v2[0:size], u[0:size_u]) \
1263  copyin(bc[0:4], do_comm[0:4], ieo, jeo, kappa, \
1264  Nst, Nst_pad, Nx, Ny, Nz, Nt)
1265  {
1266 
1267 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1268  {
1269 
1270 #pragma acc loop gang worker vector
1271  for(int ist = 0; ist < Nst; ++ist){
1272  int ix = ist % Nx;
1273  int iy = (ist/Nx) % Ny;
1274  int iz = (ist/(Nx*Ny)) % Nz;
1275  int it = ist/(Nx*Ny*Nz);
1276  int keo = (jeo + iy + iz + it) % 2;
1277 
1278  real_t v2L[NVC*ND];
1279  for(int id = 0; id < ND; ++id){
1280  for(int ivc = 0; ivc < NVC; ++ivc){
1281  v2L[ivc + NVC * id] = 0.0;
1282  }
1283  }
1284  int opr_any = 0;
1285 
1286 #include "inc/mult_Wilson_eo_2xyz_openacc-inc.h"
1287 
1288  // idir = 3
1289  if(do_comm[3] > 0){
1290 
1291  if(it == Nt-1){
1292  int ixyz = ix + Nx * (iy + Ny * iz);
1293  int istu = ist + Nst_pad * (ieo + 2*3);
1294 
1295  real_t vt1[NVC], vt2[NVC], ut[NVC];
1296  real_t wt1[2], wt2[2];
1297 
1298  for(int ivc = 0; ivc < NVC; ++ivc){
1299  vt1[ivc] = buf_tp[IDXBF(ivc, 0, ixyz)];
1300  vt2[ivc] = buf_tp[IDXBF(ivc, 1, ixyz)];
1301  }
1302 
1303  for(int ic = 0; ic < NC; ++ic){
1304 
1305  for(int ic2 = 0; ic2 < NC; ++ic2){
1306  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, istu)];
1307  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, istu)];
1308  }
1309 
1310  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1311  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1312  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1313  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1314  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1315  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1316  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1317  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1318 
1319  v2L[ic*2 + ID1] += wt1[0];
1320  v2L[ic*2+1 + ID1] += wt1[1];
1321  v2L[ic*2 + ID2] += wt2[0];
1322  v2L[ic*2+1 + ID2] += wt2[1];
1323  v2L[ic*2 + ID3] += wt1[0];
1324  v2L[ic*2+1 + ID3] += wt1[1];
1325  v2L[ic*2 + ID4] += wt2[0];
1326  v2L[ic*2+1 + ID4] += wt2[1];
1327  }
1328  ++opr_any;
1329  }
1330 
1331  if(it == 0){
1332  int ixyz = ix + Nx * (iy + Ny * iz);
1333  real_t bc2 = bc[3];
1334  real_t wt1[2], wt2[2];
1335 
1336  for(int ic = 0; ic < NC; ++ic){
1337  wt1[0] = bc2 * buf_tm[IDXBF_R(ic, 0, ixyz)];
1338  wt1[1] = bc2 * buf_tm[IDXBF_I(ic, 0, ixyz)];
1339  wt2[0] = bc2 * buf_tm[IDXBF_R(ic, 1, ixyz)];
1340  wt2[1] = bc2 * buf_tm[IDXBF_I(ic, 1, ixyz)];
1341  v2L[ic*2 + ID1] += wt1[0];
1342  v2L[ic*2+1 + ID1] += wt1[1];
1343  v2L[ic*2 + ID2] += wt2[0];
1344  v2L[ic*2+1 + ID2] += wt2[1];
1345  v2L[ic*2 + ID3] -= wt1[0];
1346  v2L[ic*2+1 + ID3] -= wt1[1];
1347  v2L[ic*2 + ID4] -= wt2[0];
1348  v2L[ic*2+1 + ID4] -= wt2[1];
1349  }
1350  ++opr_any;
1351  }
1352 
1353  }
1354 
1355  if(opr_any > 0){
1356  if(iflag == 0){
1357  for(int id = 0; id < ND; ++id){
1358  for(int ivc = 0; ivc < NVC; ++ivc){
1359  v2[IDX2_SP(ivc, id, ist)] += -kappa * v2L[ivc + NVC * id];
1360  }
1361  }
1362  }else{
1363  for(int id = 0; id < ND; ++id){
1364  for(int ivc = 0; ivc < NVC; ++ivc){
1365  v2[IDX2_SP(ivc, id, ist)] += kappa * v2L[ivc + NVC * id];
1366  }
1367  }
1368  }
1369  }
1370 
1371  }
1372 
1373  } // acc parallel
1374  } // acc data
1375 
1376 }
1377 
1378 //============================================================END=====
u14
u14
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:100
v2_24
v2_24
Definition: mult_Clover_csw_chiral_openacc-inc.h:105
v2_22
v2_22
Definition: mult_Clover_csw_chiral_openacc-inc.h:91
u13
u13
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:99
mult_wilson_eo_1_dirac
void mult_wilson_eo_1_dirac(real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm, real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym, real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm, real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm, real_t *RESTRICT u, real_t *RESTRICT v1, const int ieo, const int jeo, int *Nsize, int *bc, int *do_comm, int Nc)
Definition: mult_Wilson_eo_openacc-inc.h:261
v2_53
v2_53
Definition: mult_Clover_csw_chiral_openacc-inc.h:101
v2_01
v2_01
Definition: mult_Clover_csw_chiral_openacc-inc.h:82
mult_wilson_Meo_dirac
void mult_wilson_Meo_dirac(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, real_t *RESTRICT x1, const int ieo, const int jeo, real_t kappa, int *Nsize, int *bc, int iflag)
Definition: mult_Wilson_eo_openacc-inc.h:20
u11
u11
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:74
v2_52
v2_52
Definition: mult_Clover_csw_chiral_openacc-inc.h:94
v2_14
v2_14
Definition: mult_Clover_csw_chiral_openacc-inc.h:104
vt2_1
vt2_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:40
v2_04
v2_04
Definition: mult_Clover_csw_chiral_openacc-inc.h:103
iy
int iy
Definition: mult_Wilson_xyz_openacc-inc.h:239
ixy
int ixy
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:513
ID1
#define ID1
Definition: fopr_Wilson_impl_SU2-inc.h:18
izt
int izt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:262
u_0
u_0
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:46
v2_11
v2_11
Definition: mult_Clover_csw_chiral_openacc-inc.h:83
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
IDXBF_I
#define IDXBF_I(ic, id, ist)
Definition: define_index.h:42
vt2_5
vt2_5
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:44
u_9
u_9
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:72
vt1_2
vt1_2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:34
u_5
u_5
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:51
wt1i
wt1i
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:55
v2_13
v2_13
Definition: mult_Clover_csw_chiral_openacc-inc.h:97
vt2_2
vt2_2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:41
wt2r
wt2r
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:57
v2_23
v2_23
Definition: mult_Clover_csw_chiral_openacc-inc.h:98
u_3
u_3
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:49
mult_wilson_Meo_chiral
void mult_wilson_Meo_chiral(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, real_t *RESTRICT x1, const int ieo, const int jeo, real_t kappa, int *Nsize, int *bc, int iflag)
Definition: mult_Wilson_eo_openacc-inc.h:140
mult_wilson_eo_2_dirac
void mult_wilson_eo_2_dirac(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm, real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym, real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm, real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm, real_t kappa, const int ieo, const int jeo, const int iflag, int *Nsize, int *bc, int *do_comm, int Nc)
Definition: mult_Wilson_eo_openacc-inc.h:1062
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
IDX2_G_I
#define IDX2_G_I(ic1, ic2, ist)
Definition: define_index.h:52
v2_21
v2_21
Definition: mult_Clover_csw_chiral_openacc-inc.h:84
u_8
u_8
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:71
ID2
#define ID2
Definition: fopr_Wilson_impl_SU2-inc.h:19
ID4
#define ID4
Definition: fopr_Wilson_impl_SU2-inc.h:21
iyzt
int iyzt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:14
ix
int ix
Definition: mult_Wilson_xyz_openacc-inc.h:16
vt1_0
vt1_0
Definition: mult_Wilson_eo_t_chiral_openacc-inc.h:18
u10
u10
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:73
NC
#define NC
Definition: field_F_imp_SU2-inc.h:15
vt1
real_t vt1
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
bc2
int bc2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:62
wt1r
wt1r
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:53
wt2i
wt2i
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:59
mult_wilson_eo_2_chiral
void mult_wilson_eo_2_chiral(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm, real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym, real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm, real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm, real_t kappa, const int ieo, const int jeo, const int iflag, int *Nsize, int *bc, int *do_comm, int Nc)
Definition: mult_Wilson_eo_openacc-inc.h:1216
mult_Wilson_eo_xyz_openacc-inc.h
IDXBF
#define IDXBF(ivc, id, ist)
Definition: define_index.h:43
v2_31
v2_31
Definition: mult_Clover_csw_chiral_openacc-inc.h:85
u_2
u_2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:48
u16
u16
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:102
u_1
u_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:47
u_4
u_4
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:50
v2_32
v2_32
Definition: mult_Clover_csw_chiral_openacc-inc.h:92
ND
#define ND
Definition: field_F_imp_SU2-inc.h:18
u17
u17
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:103
it
int it
Definition: mult_Wilson_xyz_openacc-inc.h:461
v2_42
v2_42
Definition: mult_Clover_csw_chiral_openacc-inc.h:93
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
v2_03
v2_03
Definition: mult_Clover_csw_chiral_openacc-inc.h:96
mult_wilson_eo_1_chiral
void mult_wilson_eo_1_chiral(real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm, real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym, real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm, real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm, real_t *RESTRICT u, real_t *RESTRICT v1, const int ieo, const int jeo, int *Nsize, int *bc, int *do_comm, int Nc)
Definition: mult_Wilson_eo_openacc-inc.h:663
v2_51
v2_51
Definition: mult_Clover_csw_chiral_openacc-inc.h:87
ID3
#define ID3
Definition: fopr_Wilson_impl_SU2-inc.h:20
vt1_5
vt1_5
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:37
v2_12
v2_12
Definition: mult_Clover_csw_chiral_openacc-inc.h:90
u12
u12
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:98
v2_02
v2_02
Definition: mult_Clover_csw_chiral_openacc-inc.h:89
vt2_3
vt2_3
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:42
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
v2_34
v2_34
Definition: mult_Clover_csw_chiral_openacc-inc.h:106
v2_41
v2_41
Definition: mult_Clover_csw_chiral_openacc-inc.h:86
vt1_3
vt1_3
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:35
MULT_UV_I
#define MULT_UV_I(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_Wilson_inline_openacc-inc.h:14
mult_Wilson_eo_2xyz_openacc-inc.h
v2_33
v2_33
Definition: mult_Clover_csw_chiral_openacc-inc.h:99
IDXBF_R
#define IDXBF_R(ic, id, ist)
Definition: define_index.h:41
v2_44
v2_44
Definition: mult_Clover_csw_chiral_openacc-inc.h:107
MULT_UV_R
#define MULT_UV_R(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_Wilson_inline_openacc-inc.h:13
IDX2_SP_R
#define IDX2_SP_R(ic, id, ist)
Definition: define_index.h:31
iz
int iz
Definition: mult_Wilson_xyz_openacc-inc.h:462
IDX2_SP
#define IDX2_SP(ivc, id, ist)
Definition: define_index.h:33
mult_Wilson_eo_t_chiral_openacc-inc.h
NDIM
#define NDIM
Definition: contract_4spinor.cpp:18
ixyz
int ixyz
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:13
IDX2_G_R
#define IDX2_G_R(ic1, ic2, ist)
Definition: define_index.h:51
u15
u15
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:101
vt2_4
vt2_4
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:43
u_6
u_6
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:69
v2_43
v2_43
Definition: mult_Clover_csw_chiral_openacc-inc.h:100
vt2_0
vt2_0
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:39
vt2
real_t vt2
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
ND2
#define ND2
Definition: define_params_SU3.h:18
IDX2_SP_I
#define IDX2_SP_I(ic, id, ist)
Definition: define_index.h:32
u_7
u_7
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:70
mult_Wilson_eo_t_dirac_openacc-inc.h
v2_54
v2_54
Definition: mult_Clover_csw_chiral_openacc-inc.h:108
vt1_4
vt1_4
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:36
vt1_1
vt1_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:33