Bridge++  Ver.2.1.3
mult_Wilson_dir_openacc-inc.h
Go to the documentation of this file.
1 
10 #define MULT_UV_R(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_UV_I(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 
13 //====================================================================
14 void mult_wilson_xp1(real_t *RESTRICT buf, real_t *RESTRICT v1,
15  int *Nsize, int *bc, int Nc)
16 {
17  // int idir = 0;
18 
19  int Nx = Nsize[0];
20  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
21 
22  real_t bc2 = bc[0];
23 
24  int size = 2 * NC * ND * CEIL_NWP(Nx * Nyzt);
25  int size_b = 2 * NC * ND2 * CEIL_NWP(Nyzt);
26 
27 #pragma acc data present(v1[0:size], buf[0:size_b]), copyin(bc2, Nx, Nyzt)
28 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
29  {
30 
31 #pragma acc loop gang worker vector
32  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
33  int ix = 0;
34  int ist = ix + Nx * iyzt;
35 
36  real_t vt[NVC*ND];
37 
38  for(int id = 0; id < ND; ++id){
39  for(int ic = 0; ic < NC; ++ic){
40  vt[2*ic + NVC*id] = v1[IDX2_SP_R(ic, id, ist)];
41  vt[2*ic+1 + NVC*id] = v1[IDX2_SP_I(ic, id, ist)];
42  }
43  }
44 
45  real_t *vt1 = &buf[ NVC * 2 * iyzt];
46  real_t *vt2 = &buf[NVC + NVC * 2 * iyzt];
47 
48  for(int ic = 0; ic < NC; ++ic){
49  int icr = 2*ic;
50  int ici = 2*ic + 1;
51  vt1[icr] = bc2 *(vt[icr + ID1] - vt[ici + ID4]);
52  vt1[ici] = bc2 *(vt[ici + ID1] + vt[icr + ID4]);
53  vt2[icr] = bc2 *(vt[icr + ID2] - vt[ici + ID3]);
54  vt2[ici] = bc2 *(vt[ici + ID2] + vt[icr + ID3]);
55  }
56 
57  }
58 
59  }
60 
61 }
62 
63 //====================================================================
64 void mult_wilson_xp2(real_t *RESTRICT v2, real_t *RESTRICT u,
65  real_t *RESTRICT buf,
66  int *Nsize, int *bc, int Nc)
67 {
68  // int idir = 0;
69 
70  int Nx = Nsize[0];
71  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
72  int Nst = Nx * Nyzt;
73 
74  int size = NVC * ND * CEIL_NWP(Nst);
75  int size_b = NVC * ND2 * CEIL_NWP(Nyzt);
76  int size_u = NDF * CEIL_NWP(Nst);
77 
78 #pragma acc data present(v2[0:size], buf[0:size_b], u[0:size_u]), \
79  copyin(Nx, Nyzt)
80 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
81  {
82 
83 #pragma acc loop gang worker vector
84  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
85  int ix = Nx-1;
86  int ist = ix + Nx * iyzt;
87 
88  real_t vt1[NVC], vt2[NVC], ut[NDF];
89  real_t wt1[2], wt2[2];
90 
91  for(int ivc = 0; ivc < NVC; ++ivc){
92  vt1[ivc] = buf[ivc + 2 * NVC * iyzt];
93  vt2[ivc] = buf[ivc + NVC + 2 * NVC * iyzt];
94  }
95 
96  for(int ic = 0; ic < NC; ++ic){
97 
98  for(int ic2 = 0; ic2 < NC; ++ic2){
99  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
100  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
101  }
102 
103  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
104  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
105  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
106  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
107  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
108  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
109  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
110  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
111 
112  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
113  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
114  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
115  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
116  v2[IDX2_SP_R(ic, 2, ist)] += wt2[1];
117  v2[IDX2_SP_I(ic, 2, ist)] += -wt2[0];
118  v2[IDX2_SP_R(ic, 3, ist)] += wt1[1];
119  v2[IDX2_SP_I(ic, 3, ist)] += -wt1[0];
120 
121  }
122 
123  }
124 
125  }
126 
127 }
128 
129 //====================================================================
130 void mult_wilson_xpb(real_t *RESTRICT v2, real_t *RESTRICT u,
131  real_t *RESTRICT v1,
132  int *Nsize, int *bc, int Nc)
133 {
134  // int idir = 0;
135 
136  int Nx = Nsize[0];
137  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
138  int Nst = Nx * Nyzt;
139 
140  int size = NVC * ND * CEIL_NWP(Nst);
141  int size_u = NDF * CEIL_NWP(Nst);
142 
143 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
144  copyin(bc[0:4], Nx, Nyzt, Nst)
145 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
146  {
147  for(int ist = 0; ist < Nst; ++ist){
148  int ix = ist % Nx;
149  int iyzt = ist/Nx;
150  int nei = ((ix+1) % Nx) + Nx * iyzt;
151  real_t bc2 = 1.0;
152  if(ix == Nx-1) bc2 = bc[0];
153 
154  real_t vt[NVCD], vt1[NVC], vt2[NVC], ut[NDF];
155  real_t wt1[2], wt2[2];
156 
157  for(int id = 0; id < ND; ++id){
158  for(int ic = 0; ic < NC; ++ic){
159  int icr = 2*ic;
160  int ici = 2*ic + 1;
161  vt[icr + NVC*id] = v1[IDX2_SP_R(ic, id, nei)];
162  vt[ici + NVC*id] = v1[IDX2_SP_I(ic, id, nei)];
163  }
164  }
165 
166  for(int ic = 0; ic < NC; ++ic){
167  int icr = 2*ic;
168  int ici = 2*ic + 1;
169  vt1[icr] = bc2 * (vt[icr + ID1] - vt[ici + ID4]);
170  vt1[ici] = bc2 * (vt[ici + ID1] + vt[icr + ID4]);
171  vt2[icr] = bc2 * (vt[icr + ID2] - vt[ici + ID3]);
172  vt2[ici] = bc2 * (vt[ici + ID2] + vt[icr + ID3]);
173  }
174 
175  for(int ic = 0; ic < NC; ++ic){
176 
177  for(int ic2 = 0; ic2 < NC; ++ic2){
178  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
179  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
180  }
181 
182  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
183  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
184  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
185  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
186  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
187  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
188  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
189  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
190 
191  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
192  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
193  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
194  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
195  v2[IDX2_SP_R(ic, 2, ist)] += wt2[1];
196  v2[IDX2_SP_I(ic, 2, ist)] += -wt2[0];
197  v2[IDX2_SP_R(ic, 3, ist)] += wt1[1];
198  v2[IDX2_SP_I(ic, 3, ist)] += -wt1[0];
199 
200  }
201 
202  }
203 
204  }
205 
206 }
207 
208 //====================================================================
209 void mult_wilson_xm1(real_t *RESTRICT buf, real_t *RESTRICT u,
210  real_t *RESTRICT v1,
211  int *Nsize, int *bc, int Nc)
212 {
213  // int idir = 0;
214 
215  int Nx = Nsize[0];
216  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
217  int Nst = Nx * Nyzt;
218 
219  int size = NVC * ND * CEIL_NWP(Nst);
220  int size_b = NVC * ND2 * CEIL_NWP(Nyzt);
221  int size_u = NDF * CEIL_NWP(Nst);
222 
223 #pragma acc data present(v1[0:size], buf[0:size_b], u[0:size_u]), \
224  copyin(Nx, Nyzt)
225 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
226  {
227 
228 #pragma acc loop gang worker vector
229  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
230  int ix = Nx-1;
231  int ist = ix + Nx * iyzt;
232 
233  real_t vt[NVC*ND], vt1[NVC], vt2[NVC], ut[NDF];
234 
235 
236  for(int id = 0; id < ND; ++id){
237  for(int ic = 0; ic < NC; ++ic){
238  vt[2*ic + NVC * id] = v1[IDX2_SP_R(ic, id, ist)];
239  vt[2*ic+1 + NVC * id] = v1[IDX2_SP_I(ic, id, ist)];
240  }
241  }
242 
243  for(int ic = 0; ic < NC; ++ic){
244  int icr = 2*ic;
245  int ici = 2*ic + 1;
246  vt1[icr] = vt[icr + ID1] + vt[ici + ID4];
247  vt1[ici] = vt[ici + ID1] - vt[icr + ID4];
248  vt2[icr] = vt[icr + ID2] + vt[ici + ID3];
249  vt2[ici] = vt[ici + ID2] - vt[icr + ID3];
250  }
251 
252  real_t *wt1 = &buf[ NVC * 2 * iyzt];
253  real_t *wt2 = &buf[NVC + NVC * 2 * iyzt];
254 
255  for(int ic = 0; ic < NC; ++ic){
256  int icr = 2*ic;
257  int ici = 2*ic + 1;
258 
259  for(int ic2 = 0; ic2 < NC; ++ic2){
260  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, ist)];
261  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, ist)];
262  }
263 
264  wt1[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
265  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
266  wt1[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
267  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
268  wt2[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
269  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
270  wt2[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
271  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
272 
273  }
274 
275  }
276 
277  }
278 
279 }
280 
281 //====================================================================
282 void mult_wilson_xm2(real_t *RESTRICT v2, real_t *RESTRICT buf,
283  int *Nsize, int *bc, int Nc)
284 {
285  // int idir = 0;
286 
287  int Nx = Nsize[0];
288  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
289  int Nst = Nx * Nyzt;
290 
291  real_t bc2 = bc[0];
292 
293  int size = 2 * NC * ND * CEIL_NWP(Nst);
294  int size_b = 2 * NC * ND2 * CEIL_NWP(Nyzt);
295 
296 #pragma acc data present(v2[0:size], buf[0:size_b]), copyin(bc2, Nx, Nyzt)
297 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
298  {
299 
300 #pragma acc loop gang worker vector
301  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
302  int ix = 0;
303  int ist = ix + Nx * iyzt;
304 
305  real_t wt1[2], wt2[2];
306 
307  for(int ic = 0; ic < NC; ++ic){
308  int icr = 2 * ic;
309  int ici = 2 * ic + 1;
310 
311  wt1[0] = bc2 * buf[icr + 2 * NVC * iyzt];
312  wt1[1] = bc2 * buf[ici + 2 * NVC * iyzt];
313  wt2[0] = bc2 * buf[icr + NVC + 2 * NVC * iyzt];
314  wt2[1] = bc2 * buf[ici + NVC + 2 * NVC * iyzt];
315 
316  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
317  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
318  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
319  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
320  v2[IDX2_SP_R(ic, 2, ist)] += -wt2[1];
321  v2[IDX2_SP_I(ic, 2, ist)] += wt2[0];
322  v2[IDX2_SP_R(ic, 3, ist)] += -wt1[1];
323  v2[IDX2_SP_I(ic, 3, ist)] += wt1[0];
324 
325  }
326 
327  }
328 
329  }
330 
331 }
332 
333 //====================================================================
334  void mult_wilson_xmb(real_t *RESTRICT v2, real_t *RESTRICT u,
335  real_t *RESTRICT v1,
336  int *Nsize, int *bc, int Nc)
337 {
338  // int idir = 0;
339 
340  int Nx = Nsize[0];
341  int Nyzt = Nsize[1] * Nsize[2] * Nsize[3];
342  int Nst = Nx * Nyzt;
343 
344  int size = NVC * ND * CEIL_NWP(Nst);
345  int size_u = NDF * CEIL_NWP(Nst);
346 
347 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
348  copyin(bc[0:4], Nx, Nyzt, Nst)
349 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
350  {
351  for(int ist = 0; ist < Nst; ++ist){
352  int ix = ist % Nx;
353  int iyzt = ist/Nx;
354  int nei = ix-1 + Nx * iyzt;
355  if(ix == 0) nei = Nx-1 + Nx * iyzt;
356  real_t bc2 = 1.0;
357  if(ix == 0) bc2 = bc[0];
358 
359  real_t vt[NVC*ND], vt1[NVC], vt2[NVC], ut[NDF];
360  real_t wt1[2], wt2[2];
361 
362  for(int id = 0; id < ND; ++id){
363  for(int ic = 0; ic < NC; ++ic){
364  int icr = 2*ic;
365  int ici = 2*ic + 1;
366  vt[icr + NVC*id] = v1[IDX2_SP_R(ic, id, nei)];
367  vt[ici + NVC*id] = v1[IDX2_SP_I(ic, id, nei)];
368  }
369  }
370 
371  for(int ic = 0; ic < NC; ++ic){
372  int icr = 2*ic;
373  int ici = 2*ic + 1;
374  vt1[icr] = vt[icr + ID1] + vt[ici + ID4];
375  vt1[ici] = vt[ici + ID1] - vt[icr + ID4];
376  vt2[icr] = vt[icr + ID2] + vt[ici + ID3];
377  vt2[ici] = vt[ici + ID2] - vt[icr + ID3];
378  }
379 
380  for(int ic = 0; ic < NC; ++ic){
381 
382  for(int ic2 = 0; ic2 < NC; ++ic2){
383  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, nei)];
384  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, nei)];
385  }
386 
387  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
388  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
389  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
390  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
391  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
392  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
393  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
394  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
395 
396  v2[IDX2_SP_R(ic, 0, ist)] += bc2 * wt1[0];
397  v2[IDX2_SP_I(ic, 0, ist)] += bc2 * wt1[1];
398  v2[IDX2_SP_R(ic, 1, ist)] += bc2 * wt2[0];
399  v2[IDX2_SP_I(ic, 1, ist)] += bc2 * wt2[1];
400  v2[IDX2_SP_R(ic, 2, ist)] += -bc2 * wt2[1];
401  v2[IDX2_SP_I(ic, 2, ist)] += bc2 * wt2[0];
402  v2[IDX2_SP_R(ic, 3, ist)] += -bc2 * wt1[1];
403  v2[IDX2_SP_I(ic, 3, ist)] += bc2 * wt1[0];
404 
405  }
406 
407  }
408 
409  }
410 
411 }
412 
413 //====================================================================
414 void mult_wilson_yp1(real_t *RESTRICT buf, real_t *RESTRICT v1,
415  int *Nsize, int *bc, int Nc)
416 {
417  int Nx = Nsize[0];
418  int Ny = Nsize[1];
419  int Nzt = Nsize[2] * Nsize[3];
420 
421  real_t bc2 = bc[1];
422 
423  int size = NVC * ND * CEIL_NWP(Nx * Ny * Nzt);
424  int size_b = NVC * ND2 * CEIL_NWP(Nx * Nzt);
425 
426 #pragma acc data present(v1[0:size], buf[0:size_b]), \
427  copyin(bc2, Nx, Ny, Nzt)
428 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
429  {
430 
431 #pragma acc loop gang worker vector collapse(2)
432  for(int izt = 0; izt < Nzt; ++izt){
433  for(int ix = 0; ix < Nx; ++ix){
434  int iy = 0;
435  int ist = ix + Nx * (iy + Ny * izt);
436  int ixzt = ix + Nx * izt;
437 
438  real_t vt[NVC*ND];
439 
440  for(int id = 0; id < ND; ++id){
441  for(int ic = 0; ic < NC; ++ic){
442  vt[2*ic + NVC*id] = v1[IDX2_SP_R(ic, id, ist)];
443  vt[2*ic+1 + NVC*id] = v1[IDX2_SP_I(ic, id, ist)];
444  }
445  }
446 
447  real_t *vt1 = &buf[ NVC * 2 * ixzt];
448  real_t *vt2 = &buf[NVC + NVC * 2 * ixzt];
449 
450  for(int ic = 0; ic < NC; ++ic){
451  int icr = 2*ic;
452  int ici = 2*ic + 1;
453  vt1[icr] = bc2 * (vt[icr + ID1] + vt[icr + ID4]);
454  vt1[ici] = bc2 * (vt[ici + ID1] + vt[ici + ID4]);
455  vt2[icr] = bc2 * (vt[icr + ID2] - vt[icr + ID3]);
456  vt2[ici] = bc2 * (vt[ici + ID2] - vt[ici + ID3]);
457  }
458 
459  }
460  }
461 
462  }
463 
464 }
465 
466 //====================================================================
467 void mult_wilson_yp2(real_t *RESTRICT v2, real_t *RESTRICT u,
468  real_t *RESTRICT buf,
469  int *Nsize, int *bc, int Nc)
470 {
471  int idir = 1;
472 
473  int Nx = Nsize[0];
474  int Ny = Nsize[1];
475  int Nzt = Nsize[2] * Nsize[3];
476  int Nst = Nx * Ny * Nzt;
477 
478  int size = NVC * ND * CEIL_NWP(Nst);
479  int size_b = NVC * ND2 * CEIL_NWP(Nx * Nzt);
480  int size_u = NDF * CEIL_NWP(Nst);
481 
482 #pragma acc data present(v2[0:size], buf[0:size_b], u[0:size_u]), \
483  copyin(Nx, Ny, Nzt, Nst, idir)
484 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
485  {
486 
487 #pragma acc loop gang worker vector collapse(2)
488  for(int izt = 0; izt < Nzt; ++izt){
489  for(int ix = 0; ix < Nx; ++ix){
490  int iy = Ny-1;
491  int ist = ix + Nx * (iy + Ny * izt);
492  int ixzt = ix + Nx * izt;
493 
494  real_t vt1[NVC], vt2[NVC], ut[NDF];
495  real_t wt1[2], wt2[2];
496 
497  for(int ivc = 0; ivc < NVC; ++ivc){
498  vt1[ivc] = buf[ivc + 2 * NVC * ixzt];
499  vt2[ivc] = buf[ivc + NVC + 2 * NVC * ixzt];
500  }
501 
502  for(int ic = 0; ic < NC; ++ic){
503 
504  for(int ic2 = 0; ic2 < NC; ++ic2){
505  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
506  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
507  }
508 
509  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
510  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
511  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
512  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
513  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
514  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
515  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
516  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
517 
518  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
519  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
520  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
521  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
522  v2[IDX2_SP_R(ic, 2, ist)] += -wt2[0];
523  v2[IDX2_SP_I(ic, 2, ist)] += -wt2[1];
524  v2[IDX2_SP_R(ic, 3, ist)] += wt1[0];
525  v2[IDX2_SP_I(ic, 3, ist)] += wt1[1];
526 
527  }
528 
529  }
530  }
531 
532  }
533 
534 }
535 
536 //====================================================================
537  void mult_wilson_ypb(real_t *RESTRICT v2, real_t *RESTRICT u,
538  real_t *RESTRICT v1,
539  int *Nsize, int *bc, int Nc)
540 {
541  int idir = 1;
542 
543  int Nx = Nsize[0];
544  int Ny = Nsize[1];
545  int Nzt = Nsize[2] * Nsize[3];
546  int Nst = Nx * Ny * Nzt;
547 
548  int size = 2 * NC * ND * CEIL_NWP(Nst);
549  int size_u = NDF * CEIL_NWP(Nst);
550 
551 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
552  copyin(bc[0:4], idir, Nx, Ny, Nzt, Nst)
553 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
554  {
555 
556 #pragma acc loop gang worker vector
557  for(int ist = 0; ist < Nst; ++ist){
558  int ix = ist % Nx;
559  int iy = (ist/Nx) % Ny;
560  int izt = ist/(Nx*Ny);
561  int iy2 = (iy+1) % Ny;
562  int nei = ix + Nx * (iy2 + Ny * izt);
563  real_t bc2 = 1.0;
564  if(iy == Ny-1) bc2 = bc[idir];
565 
566  real_t vt[NVC*ND], vt1[NVC], vt2[NVC], ut[NDF];
567  real_t wt1[2], wt2[2];
568 
569 
570  for(int id = 0; id < ND; ++id){
571  for(int ic = 0; ic < NC; ++ic){
572  int icr = 2*ic;
573  int ici = 2*ic + 1;
574  vt[icr + NVC*id] = v1[IDX2_SP_R(ic, id, nei)];
575  vt[ici + NVC*id] = v1[IDX2_SP_I(ic, id, nei)];
576  }
577  }
578 
579  for(int ic = 0; ic < NC; ++ic){
580  int icr = 2*ic;
581  int ici = 2*ic + 1;
582  vt1[icr] = bc2 * (vt[icr + ID1] + vt[icr + ID4]);
583  vt1[ici] = bc2 * (vt[ici + ID1] + vt[ici + ID4]);
584  vt2[icr] = bc2 * (vt[icr + ID2] - vt[icr + ID3]);
585  vt2[ici] = bc2 * (vt[ici + ID2] - vt[ici + ID3]);
586  }
587 
588  for(int ic = 0; ic < NC; ++ic){
589 
590  for(int ic2 = 0; ic2 < NC; ++ic2){
591  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
592  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
593  }
594 
595  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
596  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
597  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
598  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
599  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
600  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
601  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
602  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
603 
604  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
605  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
606  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
607  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
608  v2[IDX2_SP_R(ic, 2, ist)] += -wt2[0];
609  v2[IDX2_SP_I(ic, 2, ist)] += -wt2[1];
610  v2[IDX2_SP_R(ic, 3, ist)] += wt1[0];
611  v2[IDX2_SP_I(ic, 3, ist)] += wt1[1];
612  }
613 
614  }
615 
616  }
617 
618 }
619 
620 //====================================================================
621 void mult_wilson_ym1(real_t *RESTRICT buf, real_t *RESTRICT u,
622  real_t *RESTRICT v1,
623  int *Nsize, int *bc, int Nc)
624 {
625  int idir = 1;
626 
627  int Nx = Nsize[0];
628  int Ny = Nsize[1];
629  int Nzt = Nsize[2] * Nsize[3];
630  int Nst = Nx * Ny * Nzt;
631 
632  int size = NVC * ND * CEIL_NWP(Nst);
633  int size_b = NVC * ND2 * CEIL_NWP(Nx * Nzt);
634  int size_u = NDF * CEIL_NWP(Nst);
635 
636 #pragma acc data present(v1[0:size], buf[0:size_b], u[0:size_u]), \
637  copyin(Nx, Ny, Nzt, Nst, idir)
638 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
639  {
640 
641 #pragma acc loop gang worker vector collapse(2)
642  for(int izt = 0; izt < Nzt; ++izt){
643  for(int ix = 0; ix < Nx; ++ix){
644  int iy = Ny-1;
645  int ist = ix + Nx * (iy + Ny*izt);
646  int ixzt = ix + Nx * izt;
647 
648  real_t vt[NVCD], vt1[NVC], vt2[NVC], ut[NDF];
649 
650  for(int id = 0; id < ND; ++id){
651  for(int ic = 0; ic < NC; ++ic){
652  vt[2*ic + NVC*id] = v1[IDX2_SP_R(ic, id, ist)];
653  vt[2*ic+1 + NVC*id] = v1[IDX2_SP_I(ic, id, ist)];
654  }
655  }
656 
657  for(int ic = 0; ic < NC; ++ic){
658  int icr = 2*ic;
659  int ici = 2*ic + 1;
660  vt1[icr] = vt[icr + ID1] - vt[icr + ID4];
661  vt1[ici] = vt[ici + ID1] - vt[ici + ID4];
662  vt2[icr] = vt[icr + ID2] + vt[icr + ID3];
663  vt2[ici] = vt[ici + ID2] + vt[ici + ID3];
664  }
665 
666  real_t *wt1 = &buf[ NVC * 2 * ixzt];
667  real_t *wt2 = &buf[NVC + NVC * 2 * ixzt];
668 
669  for(int ic = 0; ic < NC; ++ic){
670  int icr = 2*ic;
671  int ici = 2*ic + 1;
672 
673  for(int ic2 = 0; ic2 < NC; ++ic2){
674  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, ist)];
675  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, ist)];
676  }
677 
678  wt1[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
679  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
680  wt1[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
681  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
682  wt2[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
683  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
684  wt2[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
685  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
686  }
687 
688  }
689  }
690 
691  }
692 
693 }
694 
695 //====================================================================
696 void mult_wilson_ym2(real_t *RESTRICT v2, real_t *RESTRICT buf,
697  int *Nsize, int *bc, int Nc)
698 {
699  int idir = 1;
700 
701  int Nx = Nsize[0];
702  int Ny = Nsize[1];
703  int Nzt = Nsize[2] * Nsize[3];
704  int Nst = Nx * Ny * Nzt;
705 
706  real_t bc2 = bc[idir];
707 
708  int size = NVC * ND * CEIL_NWP(Nst);
709  int size_b = NVC * ND2 * CEIL_NWP(Nx * Nzt);
710 
711 #pragma acc data present(v2[0:size], buf[0:size_b]), \
712  copyin(bc2, Nx, Ny, Nzt)
713 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
714  {
715 
716 #pragma acc loop gang worker vector collapse(2)
717  for(int izt = 0; izt < Nzt; ++izt){
718  for(int ix = 0; ix < Nx; ++ix){
719  int iy = 0;
720  int ist = ix + Nx*(iy + Ny*izt);
721  int ixzt = ix + Nx * izt;
722 
723  real_t wt1[2], wt2[2];
724 
725  for(int ic = 0; ic < NC; ++ic){
726  int icr = 2 * ic;
727  int ici = 2 * ic + 1;
728 
729  wt1[0] = bc2 * buf[icr + 2 * NVC * ixzt];
730  wt1[1] = bc2 * buf[ici + 2 * NVC * ixzt];
731  wt2[0] = bc2 * buf[icr + NVC + 2 * NVC * ixzt];
732  wt2[1] = bc2 * buf[ici + NVC + 2 * NVC * ixzt];
733 
734  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
735  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
736  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
737  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
738  v2[IDX2_SP_R(ic, 2, ist)] += wt2[0];
739  v2[IDX2_SP_I(ic, 2, ist)] += wt2[1];
740  v2[IDX2_SP_R(ic, 3, ist)] += -wt1[0];
741  v2[IDX2_SP_I(ic, 3, ist)] += -wt1[1];
742  }
743 
744  }
745  }
746 
747  }
748 
749 }
750 
751 //====================================================================
752  void mult_wilson_ymb(real_t *RESTRICT v2, real_t *RESTRICT u,
753  real_t *RESTRICT v1,
754  int *Nsize, int *bc, int Nc)
755 {
756  int idir = 1;
757 
758  int Nx = Nsize[0];
759  int Ny = Nsize[1];
760  int Nzt = Nsize[2] * Nsize[3];
761  int Nst = Nx * Ny * Nzt;
762 
763  int size = 2 * NC * ND * CEIL_NWP(Nst);
764  int size_u = NDF * CEIL_NWP(Nst);
765 
766 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
767  copyin(bc[0:4], idir, Nx, Ny, Nzt, Nst)
768 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
769  {
770 
771 #pragma acc loop gang worker vector
772  for(int ist = 0; ist < Nst; ++ist){
773  int ix = ist % Nx;
774  int iy = (ist/Nx) % Ny;
775  int izt = ist/(Nx*Ny);
776  int iy2 = (iy-1+Ny) % Ny;
777  int nei = ix + Nx * (iy2 + Ny * izt);
778 
779  real_t vt[NVC*ND], vt1[NVC], vt2[NVC], ut[NDF];
780  real_t wt1[2], wt2[2];
781 
782  real_t bc2 = 1.0;
783  if(iy == 0) bc2 = bc[idir];
784 
785  for(int id = 0; id < ND; ++id){
786  for(int ic = 0; ic < NC; ++ic){
787  int icr = 2*ic;
788  int ici = 2*ic + 1;
789  vt[icr + NVC * id] = v1[IDX2_SP_R(ic, id, nei)];
790  vt[ici + NVC * id] = v1[IDX2_SP_I(ic, id, nei)];
791  }
792  }
793 
794  for(int ic = 0; ic < NC; ++ic){
795  int icr = 2*ic;
796  int ici = 2*ic + 1;
797  vt1[icr] = vt[icr + ID1] - vt[icr + ID4];
798  vt1[ici] = vt[ici + ID1] - vt[ici + ID4];
799  vt2[icr] = vt[icr + ID2] + vt[icr + ID3];
800  vt2[ici] = vt[ici + ID2] + vt[ici + ID3];
801  }
802 
803  for(int ic = 0; ic < NC; ++ic){
804 
805  for(int ic2 = 0; ic2 < NC; ++ic2){
806  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, nei)];
807  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, nei)];
808  }
809 
810  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
811  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
812  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
813  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
814  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
815  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
816  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
817  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
818 
819  v2[IDX2_SP_R(ic, 0, ist)] += bc2 * wt1[0];
820  v2[IDX2_SP_I(ic, 0, ist)] += bc2 * wt1[1];
821  v2[IDX2_SP_R(ic, 1, ist)] += bc2 * wt2[0];
822  v2[IDX2_SP_I(ic, 1, ist)] += bc2 * wt2[1];
823  v2[IDX2_SP_R(ic, 2, ist)] += bc2 * wt2[0];
824  v2[IDX2_SP_I(ic, 2, ist)] += bc2 * wt2[1];
825  v2[IDX2_SP_R(ic, 3, ist)] += -bc2 * wt1[0];
826  v2[IDX2_SP_I(ic, 3, ist)] += -bc2 * wt1[1];
827  }
828 
829  }
830 
831  }
832 
833 }
834 
835 //====================================================================
836 void mult_wilson_zp1(real_t *RESTRICT buf, real_t *RESTRICT v1,
837  int *Nsize, int *bc, int Nc)
838 {
839  int idir = 2;
840 
841  int Nxy = Nsize[0] * Nsize[1];
842  int Nz = Nsize[2];
843  int Nt = Nsize[3];
844 
845  real_t bc2 = bc[idir];
846 
847  int size = NVC * ND * CEIL_NWP(Nxy * Nz * Nt);
848  int size_b = NVC * ND2 * CEIL_NWP(Nxy * Nt);
849 
850 #pragma acc data present(v1[0:size], buf[0:size_b]), \
851  copyin(bc2, Nxy, Nz, Nt)
852 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
853  {
854 
855 #pragma acc loop gang worker vector collapse(2)
856  for(int it = 0; it < Nt; ++it){
857  for(int ixy = 0; ixy < Nxy; ++ixy){
858  int iz = 0;
859  int ist = ixy + Nxy * (iz + Nz * it);
860  int ixyt = ixy + Nxy * it;
861 
862  real_t vt[NVC*ND];
863 
864  for(int id = 0; id < ND; ++id){
865  for(int ic = 0; ic < NC; ++ic){
866  vt[2*ic + NVC * id] = v1[IDX2_SP_R(ic, id, ist)];
867  vt[2*ic+1 + NVC * id] = v1[IDX2_SP_I(ic, id, ist)];
868  }
869  }
870 
871  real_t *vt1 = &buf[ NVC * 2 * ixyt];
872  real_t *vt2 = &buf[NVC + NVC * 2 * ixyt];
873 
874  for(int ic = 0; ic < NC; ++ic){
875  int icr = 2*ic;
876  int ici = 2*ic + 1;
877  vt1[icr] = bc2 * (vt[icr + ID1] - vt[ici + ID3]);
878  vt1[ici] = bc2 * (vt[ici + ID1] + vt[icr + ID3]);
879  vt2[icr] = bc2 * (vt[icr + ID2] + vt[ici + ID4]);
880  vt2[ici] = bc2 * (vt[ici + ID2] - vt[icr + ID4]);
881  }
882 
883  }
884  }
885 
886  }
887 
888 }
889 
890 //====================================================================
891 void mult_wilson_zp2(real_t *RESTRICT v2, real_t *RESTRICT u,
892  real_t *RESTRICT buf,
893  int *Nsize, int *bc, int Nc)
894 {
895  int idir = 2;
896 
897  int Nxy = Nsize[0] * Nsize[1];
898  int Nz = Nsize[2];
899  int Nt = Nsize[3];
900  int Nst = Nxy * Nz * Nt;
901 
902  int size = NVC * ND * CEIL_NWP(Nst);
903  int size_b = NVC * ND2 * CEIL_NWP(Nxy * Nt);
904  int size_u = NDF * CEIL_NWP(Nst);
905 
906 #pragma acc data present(v2[0:size], buf[0:size_b], u[0:size_u]), \
907  copyin(Nxy, Nz, Nt, Nst, idir)
908 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
909  {
910 
911 #pragma acc loop gang worker vector collapse(2)
912  for(int it = 0; it < Nt; ++it){
913  for(int ixy = 0; ixy < Nxy; ++ixy){
914  int iz = Nz-1;
915  int ist = ixy + Nxy * (iz + Nz * it);
916  int ixyt = ixy + Nxy * it;
917 
918  real_t vt1[NVC], vt2[NVC], ut[NDF];
919  real_t wt1[2], wt2[2];
920 
921 
922  for(int ivc = 0; ivc < NVC; ++ivc){
923  vt1[ivc] = buf[ivc + 2 * NVC * ixyt];
924  vt2[ivc] = buf[ivc + NVC + 2 * NVC * ixyt];
925  }
926 
927  for(int ic = 0; ic < NC; ++ic){
928 
929  for(int ic2 = 0; ic2 < NC; ++ic2){
930  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
931  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
932  }
933 
934  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
935  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
936  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
937  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
938  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
939  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
940  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
941  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
942 
943  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
944  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
945  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
946  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
947  v2[IDX2_SP_R(ic, 2, ist)] += wt1[1];
948  v2[IDX2_SP_I(ic, 2, ist)] += -wt1[0];
949  v2[IDX2_SP_R(ic, 3, ist)] += -wt2[1];
950  v2[IDX2_SP_I(ic, 3, ist)] += wt2[0];
951 
952  }
953 
954  }
955  }
956 
957  }
958 
959 }
960 
961 //====================================================================
962  void mult_wilson_zpb(real_t *RESTRICT v2, real_t *RESTRICT u,
963  real_t *RESTRICT v1,
964  int *Nsize, int *bc, int Nc)
965 {
966  int idir = 2;
967 
968  int Nxy = Nsize[0] * Nsize[1];
969  int Nz = Nsize[2];
970  int Nt = Nsize[3];
971  int Nst = Nxy * Nz * Nt;
972 
973  int size = NVC * ND * CEIL_NWP(Nst);
974  int size_u = NDF * CEIL_NWP(Nst);
975 
976 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
977  copyin(bc[0:4], idir, Nxy, Nz, Nt, Nst)
978 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
979  {
980 
981 #pragma acc loop gang worker vector
982  for(int ist = 0; ist < Nst; ++ist){
983  int ixy = ist % Nxy;
984  int iz = (ist/Nxy) % Nz;
985  int it = ist/(Nxy*Nz);
986  int iz2 = (iz+1) % Nz;
987  int nei = ixy + Nxy * (iz2 + Nz * it);
988  real_t bc2 = 1.0;
989  if(iz == Nz-1) bc2 = bc[idir];
990 
991  real_t vt[NVC*ND], vt1[NVC], vt2[NVC], ut[NDF];
992  real_t wt1[2], wt2[2];
993 
994  for(int id = 0; id < ND; ++id){
995  for(int ic = 0; ic < NC; ++ic){
996  int icr = 2*ic;
997  int ici = 2*ic + 1;
998  vt[icr + NVC*id] = v1[IDX2_SP_R(ic, id, nei)];
999  vt[ici + NVC*id] = v1[IDX2_SP_I(ic, id, nei)];
1000  }
1001  }
1002 
1003  for(int ic = 0; ic < NC; ++ic){
1004  int icr = 2*ic;
1005  int ici = 2*ic + 1;
1006  vt1[icr] = bc2 * (vt[icr + ID1] - vt[ici + ID3]);
1007  vt1[ici] = bc2 * (vt[ici + ID1] + vt[icr + ID3]);
1008  vt2[icr] = bc2 * (vt[icr + ID2] + vt[ici + ID4]);
1009  vt2[ici] = bc2 * (vt[ici + ID2] - vt[icr + ID4]);
1010  }
1011 
1012  for(int ic = 0; ic < NC; ++ic){
1013 
1014  for(int ic2 = 0; ic2 < NC; ++ic2){
1015  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
1016  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
1017  }
1018 
1019  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1020  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1021  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1022  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1023  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1024  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1025  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1026  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1027 
1028  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
1029  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
1030  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
1031  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
1032  v2[IDX2_SP_R(ic, 2, ist)] += wt1[1];
1033  v2[IDX2_SP_I(ic, 2, ist)] += -wt1[0];
1034  v2[IDX2_SP_R(ic, 3, ist)] += -wt2[1];
1035  v2[IDX2_SP_I(ic, 3, ist)] += wt2[0];
1036  }
1037 
1038  }
1039 
1040  }
1041 
1042 }
1043 
1044 //====================================================================
1045 void mult_wilson_zm1(real_t *RESTRICT buf, real_t *RESTRICT u,
1046  real_t *RESTRICT v1,
1047  int *Nsize, int *bc, int Nc)
1048 {
1049  int idir = 2;
1050 
1051  int Nxy = Nsize[0] * Nsize[1];
1052  int Nz = Nsize[2];
1053  int Nt = Nsize[3];
1054  int Nst = Nxy * Nz * Nt;
1055 
1056  int size = NVC * ND * CEIL_NWP(Nst);
1057  int size_b = NVC * ND2 * CEIL_NWP(Nxy * Nt);
1058  int size_u = NDF * CEIL_NWP(Nst);
1059 
1060 #pragma acc data present(v1[0:size], buf[0:size_b], u[0:size_u]), \
1061  copyin(Nxy, Nz, Nt, Nst, idir)
1062 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1063  {
1064 
1065 #pragma acc loop gang worker vector collapse(2)
1066  for(int it = 0; it < Nt; ++it){
1067  for(int ixy = 0; ixy < Nxy; ++ixy){
1068  int iz = Nz-1;
1069  int ist = ixy + Nxy * (iz + Nz * it);
1070  int ixyt = ixy + Nxy * it;
1071 
1072  real_t vt[NVCD], vt1[NVC], vt2[NVC], ut[NDF];
1073 
1074  for(int id = 0; id < ND; ++id){
1075  for(int ic = 0; ic < NC; ++ic){
1076  vt[2*ic + NVC*id] = v1[IDX2_SP_R(ic, id, ist)];
1077  vt[2*ic+1 + NVC*id] = v1[IDX2_SP_I(ic, id, ist)];
1078  }
1079  }
1080 
1081  for(int ic = 0; ic < NC; ++ic){
1082  int icr = 2*ic;
1083  int ici = 2*ic + 1;
1084  vt1[icr] = vt[icr + ID1] + vt[ici + ID3];
1085  vt1[ici] = vt[ici + ID1] - vt[icr + ID3];
1086  vt2[icr] = vt[icr + ID2] - vt[ici + ID4];
1087  vt2[ici] = vt[ici + ID2] + vt[icr + ID4];
1088  }
1089 
1090  real_t *wt1 = &buf[ NVC * 2 * ixyt];
1091  real_t *wt2 = &buf[NVC + NVC * 2 * ixyt];
1092 
1093  for(int ic = 0; ic < NC; ++ic){
1094  int icr = 2*ic;
1095  int ici = 2*ic + 1;
1096 
1097  for(int ic2 = 0; ic2 < NC; ++ic2){
1098  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, ist)];
1099  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, ist)];
1100  }
1101 
1102  wt1[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1103  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1104  wt1[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1105  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1106  wt2[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1107  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1108  wt2[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1109  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1110 
1111  }
1112 
1113  }
1114  }
1115 
1116  }
1117 
1118 }
1119 
1120 //====================================================================
1121 void mult_wilson_zm2(real_t *RESTRICT v2, real_t *RESTRICT buf,
1122  int *Nsize, int *bc, int Nc)
1123 {
1124  int idir = 2;
1125 
1126  int Nxy = Nsize[0] * Nsize[1];
1127  int Nz = Nsize[2];
1128  int Nt = Nsize[3];
1129  int Nst = Nxy * Nz * Nt;
1130 
1131  real_t bc2 = bc[idir];
1132 
1133  int size = NVC * ND * CEIL_NWP(Nst);
1134  int size_b = NVC * ND2 * CEIL_NWP(Nxy * Nt);
1135 
1136 #pragma acc data present(v2[0:size], buf[0:size_b]), \
1137  copyin(bc2, Nxy, Nz, Nt)
1138 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1139  {
1140 
1141 #pragma acc loop gang worker vector collapse(2)
1142  for(int it = 0; it < Nt; ++it){
1143  for(int ixy = 0; ixy < Nxy; ++ixy){
1144  int iz = 0;
1145  int ist = ixy + Nxy * (iz + Nz * it);
1146  int ixyt = ixy + Nxy * it;
1147 
1148  real_t wt1[2], wt2[2];
1149 
1150  for(int ic = 0; ic < NC; ++ic){
1151  int icr = 2 * ic;
1152  int ici = 2 * ic + 1;
1153 
1154  wt1[0] = bc2 * buf[icr + 2 * NVC * ixyt];
1155  wt1[1] = bc2 * buf[ici + 2 * NVC * ixyt];
1156  wt2[0] = bc2 * buf[icr + NVC + 2 * NVC * ixyt];
1157  wt2[1] = bc2 * buf[ici + NVC + 2 * NVC * ixyt];
1158 
1159  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
1160  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
1161  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
1162  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
1163  v2[IDX2_SP_R(ic, 2, ist)] += -wt1[1];
1164  v2[IDX2_SP_I(ic, 2, ist)] += wt1[0];
1165  v2[IDX2_SP_R(ic, 3, ist)] += wt2[1];
1166  v2[IDX2_SP_I(ic, 3, ist)] += -wt2[0];
1167  }
1168 
1169  }
1170  }
1171 
1172  }
1173 
1174 }
1175 
1176 //====================================================================
1177 void mult_wilson_zmb(real_t *RESTRICT v2, real_t *RESTRICT u,
1178  real_t *RESTRICT v1,
1179  int *Nsize, int *bc, int Nc)
1180 {
1181  int idir = 2;
1182 
1183  int Nxy = Nsize[0] * Nsize[1];
1184  int Nz = Nsize[2];
1185  int Nt = Nsize[3];
1186  int Nst = Nxy * Nz * Nt;
1187 
1188  int size = 2 * NC * ND * CEIL_NWP(Nst);
1189  int size_u = NDF * CEIL_NWP(Nst);
1190 
1191 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
1192  copyin(bc[0:4], idir, Nxy, Nz, Nt, Nst)
1193 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1194  {
1195 
1196 #pragma acc loop gang worker vector
1197  for(int ist = 0; ist < Nst; ++ist){
1198  int ixy = ist % Nxy;
1199  int iz = (ist/Nxy) % Nz;
1200  int it = ist/(Nxy*Nz);
1201  int iz2 = (iz-1+Nz) % Nz;
1202  int nei = ixy + Nxy * (iz2 + Nz * it);
1203  real_t bc2 = 1.0;
1204  if(iz == 0) bc2 = bc[idir];
1205 
1206  real_t vt[NVCD], vt1[NVC], vt2[NVC], ut[NDF];
1207  real_t wt1[2], wt2[2];
1208 
1209 
1210  for(int id = 0; id < ND; ++id){
1211  for(int ic = 0; ic < NC; ++ic){
1212  int icr = 2*ic;
1213  int ici = 2*ic + 1;
1214  vt[icr + NVC*id] = v1[IDX2_SP_R(ic, id, nei)];
1215  vt[ici + NVC*id] = v1[IDX2_SP_I(ic, id, nei)];
1216  }
1217  }
1218 
1219  for(int ic = 0; ic < NC; ++ic){
1220  int icr = 2*ic;
1221  int ici = 2*ic + 1;
1222  vt1[icr] = vt[icr + ID1] + vt[ici + ID3];
1223  vt1[ici] = vt[ici + ID1] - vt[icr + ID3];
1224  vt2[icr] = vt[icr + ID2] - vt[ici + ID4];
1225  vt2[ici] = vt[ici + ID2] + vt[icr + ID4];
1226  }
1227 
1228  for(int ic = 0; ic < NC; ++ic){
1229 
1230  for(int ic2 = 0; ic2 < NC; ++ic2){
1231  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, nei)];
1232  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, nei)];
1233  }
1234 
1235  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1236  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1237  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1238  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1239  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1240  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1241  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1242  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1243 
1244  v2[IDX2_SP_R(ic, 0, ist)] += bc2 * wt1[0];
1245  v2[IDX2_SP_I(ic, 0, ist)] += bc2 * wt1[1];
1246  v2[IDX2_SP_R(ic, 1, ist)] += bc2 * wt2[0];
1247  v2[IDX2_SP_I(ic, 1, ist)] += bc2 * wt2[1];
1248  v2[IDX2_SP_R(ic, 2, ist)] += -bc2 * wt1[1];
1249  v2[IDX2_SP_I(ic, 2, ist)] += bc2 * wt1[0];
1250  v2[IDX2_SP_R(ic, 3, ist)] += bc2 * wt2[1];
1251  v2[IDX2_SP_I(ic, 3, ist)] += -bc2 * wt2[0];
1252  }
1253 
1254  }
1255 
1256  }
1257 
1258 }
1259 
1260 //====================================================================
1261 void mult_wilson_tp1_dirac(real_t *RESTRICT buf, real_t *RESTRICT v1,
1262  int *Nsize, int *bc, int Nc)
1263 {
1264  int idir = 3;
1265 
1266  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1267  int Nt = Nsize[3];
1268 
1269  real_t bc2 = bc[idir];
1270 
1271  int size = 2 * NC * ND * CEIL_NWP(Nxyz * Nt);
1272  int size_b = 2 * NC * ND2 * CEIL_NWP(Nxyz);
1273 
1274 #pragma acc data present(v1[0:size], buf[0:size_b]), copyin(bc2, Nxyz)
1275 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1276  {
1277 
1278 #pragma acc loop gang worker vector
1279  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
1280  int it = 0;
1281  int ist = ixyz + Nxyz * it;
1282 
1283  real_t vt[NVC*ND];
1284 
1285  for(int id = 2; id < ND; ++id){
1286  for(int ic = 0; ic < NC; ++ic){
1287  vt[2*ic + NVC*id] = v1[IDX2_SP_R(ic, id, ist)];
1288  vt[2*ic+1 + NVC*id] = v1[IDX2_SP_I(ic, id, ist)];
1289  }
1290  }
1291 
1292  real_t *vt1 = &buf[ NVC * 2 * ixyz];
1293  real_t *vt2 = &buf[NVC + NVC * 2 * ixyz];
1294 
1295  for(int ic = 0; ic < NC; ++ic){
1296  int icr = 2*ic;
1297  int ici = 2*ic + 1;
1298  vt1[icr] = bc2 * 2.0 * vt[icr + ID3];
1299  vt1[ici] = bc2 * 2.0 * vt[ici + ID3];
1300  vt2[icr] = bc2 * 2.0 * vt[icr + ID4];
1301  vt2[ici] = bc2 * 2.0 * vt[ici + ID4];
1302  }
1303 
1304  }
1305 
1306  }
1307 
1308 }
1309 
1310 //====================================================================
1311 void mult_wilson_tp2_dirac(real_t *RESTRICT v2, real_t *RESTRICT u,
1312  real_t *RESTRICT buf,
1313  int *Nsize, int *bc, int Nc)
1314 {
1315  int idir = 3;
1316 
1317  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1318  int Nt = Nsize[3];
1319  int Nst = Nxyz * Nt;
1320 
1321  int size = 2 * NC * ND * CEIL_NWP(Nst);
1322  int size_b = 2 * NC * ND2 * CEIL_NWP(Nxyz);
1323  int size_u = NDF * CEIL_NWP(Nst);
1324 
1325 #pragma acc data present(v2[0:size], buf[0:size_b], u[0:size_u]), \
1326  copyin(Nxyz, Nst, idir)
1327 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1328  {
1329 
1330 #pragma acc loop gang worker vector
1331  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
1332  int it = Nt-1;
1333  int ist = ixyz + Nxyz * it;
1334 
1335  real_t vt1[NVC], vt2[NVC], ut[NDF];
1336  real_t wt1[2], wt2[2];
1337 
1338  for(int ivc = 0; ivc < NVC; ++ivc){
1339  vt1[ivc] = buf[ivc + 2 * NVC * ixyz];
1340  vt2[ivc] = buf[ivc + NVC + 2 * NVC * ixyz];
1341  }
1342 
1343  for(int ic = 0; ic < NC; ++ic){
1344 
1345  for(int ic2 = 0; ic2 < NC; ++ic2){
1346  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
1347  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
1348  }
1349 
1350  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1351  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1352  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1353  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1354  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1355  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1356  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1357  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1358 
1359  v2[IDX2_SP_R(ic, 2, ist)] += wt1[0];
1360  v2[IDX2_SP_I(ic, 2, ist)] += wt1[1];
1361  v2[IDX2_SP_R(ic, 3, ist)] += wt2[0];
1362  v2[IDX2_SP_I(ic, 3, ist)] += wt2[1];
1363 
1364  }
1365 
1366  }
1367 
1368  }
1369 
1370 }
1371 
1372 //====================================================================
1373  void mult_wilson_tpb_dirac(real_t *RESTRICT v2, real_t *RESTRICT u,
1374  real_t *RESTRICT v1,
1375  int *Nsize, int *bc, int Nc)
1376 {
1377  int idir = 3;
1378 
1379  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1380  int Nt = Nsize[3];
1381  int Nst = Nxyz * Nt;
1382 
1383  int size = 2 * NC * ND * CEIL_NWP(Nst);
1384  int size_u = NDF * CEIL_NWP(Nst);
1385 
1386 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
1387  copyin(bc[0:4], idir, Nxyz, Nt, Nst)
1388 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1389  {
1390 
1391 #pragma acc loop gang worker vector
1392  for(int ist = 0; ist < Nst; ++ist){
1393  int ixyz = ist % Nxyz;
1394  int it = ist/Nxyz;
1395  int it2 = (it+1) % Nt;
1396  int nei = ixyz + Nxyz * it2;
1397  real_t bc2 = 1.0;
1398  if(it == Nt-1) bc2 = bc[idir];
1399 
1400  real_t vt[NVCD], vt1[NVC], vt2[NVC], ut[NDF];
1401  real_t wt1[2], wt2[2];
1402 
1403  for(int id = 2; id < ND; ++id){
1404  for(int ic = 0; ic < NC; ++ic){
1405  int icr = 2*ic;
1406  int ici = 2*ic + 1;
1407  vt[icr + NVC*id] = v1[IDX2_SP_R(ic, id, nei)];
1408  vt[ici + NVC*id] = v1[IDX2_SP_I(ic, id, nei)];
1409  }
1410  }
1411 
1412  for(int ic = 0; ic < NC; ++ic){
1413  int icr = 2*ic;
1414  int ici = 2*ic + 1;
1415  vt1[icr] = bc2 * 2.0 * vt[icr + ID3];
1416  vt1[ici] = bc2 * 2.0 * vt[ici + ID3];
1417  vt2[icr] = bc2 * 2.0 * vt[icr + ID4];
1418  vt2[ici] = bc2 * 2.0 * vt[ici + ID4];
1419  }
1420 
1421  for(int ic = 0; ic < NC; ++ic){
1422 
1423  for(int ic2 = 0; ic2 < NC; ++ic2){
1424  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
1425  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
1426  }
1427 
1428  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1429  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1430  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1431  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1432  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1433  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1434  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1435  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1436 
1437  v2[IDX2_SP_R(ic, 2, ist)] += wt1[0];
1438  v2[IDX2_SP_I(ic, 2, ist)] += wt1[1];
1439  v2[IDX2_SP_R(ic, 3, ist)] += wt2[0];
1440  v2[IDX2_SP_I(ic, 3, ist)] += wt2[1];
1441  }
1442 
1443  }
1444 
1445  }
1446 
1447 }
1448 
1449 //====================================================================
1450 void mult_wilson_tm1_dirac(real_t *RESTRICT buf, real_t *RESTRICT u,
1451  real_t *RESTRICT v1,
1452  int *Nsize, int *bc, int Nc)
1453 {
1454  int idir = 3;
1455 
1456  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1457  int Nt = Nsize[3];
1458  int Nst = Nxyz * Nt;
1459 
1460  int size = 2 * NC * ND * CEIL_NWP(Nst);
1461  int size_b = 2 * NC * ND2 * CEIL_NWP(Nxyz);
1462  int size_u = NDF * CEIL_NWP(Nst);
1463 
1464 #pragma acc data present(v1[0:size], buf[0:size_b], u[0:size_u]), \
1465  copyin(Nxyz, Nt, Nst, idir)
1466 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1467  {
1468 
1469 #pragma acc loop gang worker vector
1470  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
1471  int it = Nt-1;
1472  int ist = ixyz + Nxyz * it;
1473 
1474  real_t vt[NVCD], vt1[NVC], vt2[NVC], ut[NDF];
1475 
1476  for(int id = 0; id < ND2; ++id){
1477  for(int ic = 0; ic < NC; ++ic){
1478  vt[2*ic + NVC * id] = v1[IDX2_SP_R(ic, id, ist)];
1479  vt[2*ic+1 + NVC * id] = v1[IDX2_SP_I(ic, id, ist)];
1480  }
1481  }
1482 
1483  for(int ic = 0; ic < NC; ++ic){
1484  int icr = 2*ic;
1485  int ici = 2*ic + 1;
1486  vt1[icr] = 2.0 * vt[icr + ID1];
1487  vt1[ici] = 2.0 * vt[ici + ID1];
1488  vt2[icr] = 2.0 * vt[icr + ID2];
1489  vt2[ici] = 2.0 * vt[ici + ID2];
1490  }
1491 
1492  real_t *wt1 = &buf[ NVC * 2 * ixyz];
1493  real_t *wt2 = &buf[NVC + NVC * 2 * ixyz];
1494 
1495  for(int ic = 0; ic < NC; ++ic){
1496  int icr = 2*ic;
1497  int ici = 2*ic + 1;
1498 
1499  for(int ic2 = 0; ic2 < NC; ++ic2){
1500  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, ist)];
1501  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, ist)];
1502  }
1503 
1504  wt1[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1505  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1506  wt1[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1507  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1508  wt2[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1509  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1510  wt2[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1511  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1512 
1513  }
1514 
1515  }
1516 
1517  }
1518 
1519 }
1520 
1521 //====================================================================
1522 void mult_wilson_tm2_dirac(real_t *RESTRICT v2, real_t *RESTRICT buf,
1523  int *Nsize, int *bc, int Nc)
1524 {
1525  int idir = 3;
1526 
1527  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1528  int Nt = Nsize[3];
1529  int Nst = Nxyz * Nt;
1530 
1531  real_t bc2 = bc[idir];
1532 
1533  int size = 2 * NC * ND * CEIL_NWP(Nst);
1534  int size_b = 2 * NC * ND2 * CEIL_NWP(Nxyz);
1535 
1536 #pragma acc data present(v2[0:size], buf[0:size_b]), copyin(bc2, Nxyz)
1537 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1538  {
1539 
1540 #pragma acc loop gang worker vector
1541  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
1542  int it = 0;
1543  int ist = ixyz + Nxyz * it;
1544 
1545  real_t wt1[2], wt2[2];
1546 
1547  for(int ic = 0; ic < NC; ++ic){
1548  int icr = 2 * ic;
1549  int ici = 2 * ic + 1;
1550 
1551  wt1[0] = bc2 * buf[icr + 2 * NVC * ixyz];
1552  wt1[1] = bc2 * buf[ici + 2 * NVC * ixyz];
1553  wt2[0] = bc2 * buf[icr + NVC + 2 * NVC * ixyz];
1554  wt2[1] = bc2 * buf[ici + NVC + 2 * NVC * ixyz];
1555 
1556  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
1557  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
1558  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
1559  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
1560  }
1561 
1562  }
1563 
1564  }
1565 
1566 }
1567 
1568 //====================================================================
1569 void mult_wilson_tmb_dirac(real_t *RESTRICT v2, real_t *RESTRICT u,
1570  real_t *RESTRICT v1,
1571  int *Nsize, int *bc, int Nc)
1572 {
1573  int idir = 3;
1574 
1575  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1576  int Nt = Nsize[3];
1577  int Nst = Nxyz * Nt;
1578 
1579  int size = 2 * NC * ND * CEIL_NWP(Nst);
1580  int size_u = NDF * CEIL_NWP(Nst);
1581 
1582 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
1583  copyin(bc[0:4], idir, Nxyz, Nt, Nst)
1584 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1585  {
1586 
1587 #pragma acc loop gang worker vector
1588  for(int ist = 0; ist < Nst; ++ist){
1589  int ixyz = ist % Nxyz;
1590  int it = ist/Nxyz;
1591  int it2 = (it-1+Nt) % Nt;
1592  int nei = ixyz + Nxyz * it2;
1593  real_t bc2 = 1.0;
1594  if(it == 0) bc2 = bc[idir];
1595 
1596  real_t vt[NVCD];
1597  real_t vt1[NVC], vt2[NVC], ut[NDF];
1598  real_t wt1[2], wt2[2];
1599 
1600  for(int id = 0; id < ND2; ++id){
1601  for(int ic = 0; ic < NC; ++ic){
1602  int icr = 2*ic;
1603  int ici = 2*ic + 1;
1604  vt[icr + NVC*id] = v1[IDX2_SP_R(ic, id, nei)];
1605  vt[ici + NVC*id] = v1[IDX2_SP_I(ic, id, nei)];
1606  }
1607  }
1608 
1609  for(int ic = 0; ic < NC; ++ic){
1610  int icr = 2*ic;
1611  int ici = 2*ic + 1;
1612  vt1[icr] = 2.0 * vt[icr + ID1];
1613  vt1[ici] = 2.0 * vt[ici + ID1];
1614  vt2[icr] = 2.0 * vt[icr + ID2];
1615  vt2[ici] = 2.0 * vt[ici + ID2];
1616  }
1617 
1618  for(int ic = 0; ic < NC; ++ic){
1619 
1620  for(int ic2 = 0; ic2 < NC; ++ic2){
1621  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, nei)];
1622  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, nei)];
1623  }
1624 
1625  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1626  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1627  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1628  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1629  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1630  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1631  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1632  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1633 
1634  v2[IDX2_SP_R(ic, 0, ist)] += bc2 * wt1[0];
1635  v2[IDX2_SP_I(ic, 0, ist)] += bc2 * wt1[1];
1636  v2[IDX2_SP_R(ic, 1, ist)] += bc2 * wt2[0];
1637  v2[IDX2_SP_I(ic, 1, ist)] += bc2 * wt2[1];
1638  }
1639 
1640  }
1641 
1642  }
1643 
1644 }
1645 
1646 //====================================================================
1647 void mult_wilson_tp1_chiral(real_t *RESTRICT buf, real_t *RESTRICT v1,
1648  int *Nsize, int *bc, int Nc)
1649 {
1650  int idir = 3;
1651 
1652  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1653  int Nt = Nsize[3];
1654 
1655  real_t bc2 = bc[idir];
1656 
1657  int size = 2 * NC * ND * CEIL_NWP(Nxyz * Nt);
1658  int size_b = 2 * NC * ND2 * CEIL_NWP(Nxyz);
1659 
1660 #pragma acc data present(v1[0:size], buf[0:size_b]), copyin(bc2, Nxyz)
1661 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1662  {
1663 
1664 #pragma acc loop gang worker vector
1665  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
1666  int it = 0;
1667  int ist = ixyz + Nxyz * it;
1668 
1669  real_t vt[NVC*ND];
1670 
1671  for(int id = 0; id < ND; ++id){
1672  for(int ic = 0; ic < NC; ++ic){
1673  vt[2*ic + NVC * id] = v1[IDX2_SP_R(ic, id, ist)];
1674  vt[2*ic+1 + NVC * id] = v1[IDX2_SP_I(ic, id, ist)];
1675  }
1676  }
1677 
1678  real_t *vt1 = &buf[ NVC * 2 * ixyz];
1679  real_t *vt2 = &buf[NVC + NVC * 2 * ixyz];
1680 
1681  for(int ivc = 0; ivc < NVC; ++ivc){
1682  vt1[ivc] = bc2 * (vt[ivc + ID1] + vt[ivc + ID3]);
1683  vt2[ivc] = bc2 * (vt[ivc + ID2] + vt[ivc + ID4]);
1684  }
1685 
1686  }
1687 
1688  }
1689 
1690 }
1691 
1692 //====================================================================
1693 void mult_wilson_tp2_chiral(real_t *RESTRICT v2, real_t *RESTRICT u,
1694  real_t *RESTRICT buf,
1695  int *Nsize, int *bc, int Nc)
1696 {
1697  int idir = 3;
1698 
1699  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1700  int Nt = Nsize[3];
1701  int Nst = Nxyz * Nt;
1702 
1703  int size = 2 * NC * ND * CEIL_NWP(Nst);
1704  int size_b = 2 * NC * ND2 * CEIL_NWP(Nxyz);
1705  int size_u = NDF * CEIL_NWP(Nst);
1706 
1707 #pragma acc data present(v2[0:size], buf[0:size_b], u[0:size_u]), \
1708  copyin(Nxyz, Nst, idir)
1709 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1710  {
1711 
1712 #pragma acc loop gang worker vector
1713  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
1714  int it = Nt-1;
1715  int ist = ixyz + Nxyz * it;
1716 
1717  real_t vt1[NVC], vt2[NVC], ut[NDF];
1718  real_t wt1[2], wt2[2];
1719 
1720  for(int ivc = 0; ivc < NVC; ++ivc){
1721  vt1[ivc] = buf[ivc + 2 * NVC * ixyz];
1722  vt2[ivc] = buf[ivc + NVC + 2 * NVC * ixyz];
1723  }
1724 
1725  for(int ic = 0; ic < NC; ++ic){
1726 
1727  for(int ic2 = 0; ic2 < NC; ++ic2){
1728  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
1729  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
1730  }
1731 
1732  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1733  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1734  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1735  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1736  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1737  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1738  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1739  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1740 
1741  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
1742  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
1743  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
1744  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
1745  v2[IDX2_SP_R(ic, 2, ist)] += wt1[0];
1746  v2[IDX2_SP_I(ic, 2, ist)] += wt1[1];
1747  v2[IDX2_SP_R(ic, 3, ist)] += wt2[0];
1748  v2[IDX2_SP_I(ic, 3, ist)] += wt2[1];
1749 
1750  }
1751 
1752  }
1753 
1754  }
1755 
1756 }
1757 
1758 //====================================================================
1759  void mult_wilson_tpb_chiral(real_t *RESTRICT v2, real_t *RESTRICT u,
1760  real_t *RESTRICT v1,
1761  int *Nsize, int *bc, int Nc)
1762 {
1763  int idir = 3;
1764 
1765  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1766  int Nt = Nsize[3];
1767  int Nst = Nxyz * Nt;
1768 
1769  int size = 2 * NC * ND * CEIL_NWP(Nst);
1770  int size_u = NDF * CEIL_NWP(Nst);
1771 
1772 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
1773  copyin(bc[0:4], idir, Nxyz, Nt, Nst)
1774 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1775  {
1776 
1777 #pragma acc loop gang worker vector
1778  for(int ist = 0; ist < Nst; ++ist){
1779  int ixyz = ist % Nxyz;
1780  int it = ist/Nxyz;
1781  int it2 = (it+1) % Nt;
1782  int nei = ixyz + Nxyz * it2;
1783  real_t bc2 = 1.0;
1784  if(it == Nt-1) bc2 = bc[idir];
1785 
1786  real_t vt[NVCD], vt1[NVC], vt2[NVC], ut[NDF];
1787  real_t wt1[2], wt2[2];
1788 
1789 
1790  for(int id = 0; id < ND; ++id){
1791  for(int ic = 0; ic < NC; ++ic){
1792  int icr = 2*ic;
1793  int ici = 2*ic + 1;
1794  vt[icr + NVC * id] = v1[IDX2_SP_R(ic, id, nei)];
1795  vt[ici + NVC * id] = v1[IDX2_SP_I(ic, id, nei)];
1796  }
1797  }
1798 
1799  for(int ivc = 0; ivc < NVC; ++ivc){
1800  vt1[ivc] = bc2 * (vt[ivc + ID1] + vt[ivc + ID3]);
1801  vt2[ivc] = bc2 * (vt[ivc + ID2] + vt[ivc + ID4]);
1802  }
1803 
1804  for(int ic = 0; ic < NC; ++ic){
1805 
1806  for(int ic2 = 0; ic2 < NC; ++ic2){
1807  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, ist)];
1808  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, ist)];
1809  }
1810 
1811  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1812  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1813  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1814  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1815  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1816  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1817  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1818  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1819 
1820  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
1821  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
1822  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
1823  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
1824  v2[IDX2_SP_R(ic, 2, ist)] += wt1[0];
1825  v2[IDX2_SP_I(ic, 2, ist)] += wt1[1];
1826  v2[IDX2_SP_R(ic, 3, ist)] += wt2[0];
1827  v2[IDX2_SP_I(ic, 3, ist)] += wt2[1];
1828  }
1829 
1830  }
1831 
1832  }
1833 
1834 }
1835 
1836 //====================================================================
1837 void mult_wilson_tm1_chiral(real_t *RESTRICT buf, real_t *RESTRICT u,
1838  real_t *RESTRICT v1,
1839  int *Nsize, int *bc, int Nc)
1840 {
1841  int idir = 3;
1842 
1843  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1844  int Nt = Nsize[3];
1845  int Nst = Nxyz * Nt;
1846 
1847  int size = 2 * NC * ND * CEIL_NWP(Nst);
1848  int size_b = 2 * NC * ND2 * CEIL_NWP(Nxyz);
1849  int size_u = NDF * CEIL_NWP(Nst);
1850 
1851 #pragma acc data present(v1[0:size], buf[0:size_b], u[0:size_u]), \
1852  copyin(Nxyz, Nt, Nst, idir)
1853 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1854  {
1855 
1856 #pragma acc loop gang worker vector
1857  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
1858  int it = Nt-1;
1859  int ist = ixyz + Nxyz * it;
1860 
1861  real_t vt[NVCD], vt1[NVC], vt2[NVC], ut[NDF];
1862 
1863  for(int id = 0; id < ND; ++id){
1864  for(int ic = 0; ic < NC; ++ic){
1865  vt[2*ic + NVC * id] = v1[IDX2_SP_R(ic, id, ist)];
1866  vt[2*ic+1 + NVC * id] = v1[IDX2_SP_I(ic, id, ist)];
1867  }
1868  }
1869 
1870  for(int ivc = 0; ivc < NVC; ++ivc){
1871  vt1[ivc] = vt[ivc + ID1] - vt[ivc + ID3];
1872  vt2[ivc] = vt[ivc + ID2] - vt[ivc + ID4];
1873  }
1874 
1875  real_t *wt1 = &buf[ NVC * 2 * ixyz];
1876  real_t *wt2 = &buf[NVC + NVC * 2 * ixyz];
1877 
1878  for(int ic = 0; ic < NC; ++ic){
1879  int icr = 2*ic;
1880  int ici = 2*ic + 1;
1881 
1882  for(int ic2 = 0; ic2 < NC; ++ic2){
1883  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, ist)];
1884  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, ist)];
1885  }
1886 
1887  wt1[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1888  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1889  wt1[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1890  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
1891  wt2[icr] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1892  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1893  wt2[ici] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
1894  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
1895 
1896  }
1897 
1898  }
1899 
1900  }
1901 
1902 }
1903 
1904 //====================================================================
1905 void mult_wilson_tm2_chiral(real_t *RESTRICT v2, real_t *RESTRICT buf,
1906  int *Nsize, int *bc, int Nc)
1907 {
1908  int idir = 3;
1909 
1910  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1911  int Nt = Nsize[3];
1912  int Nst = Nxyz * Nt;
1913 
1914  real_t bc2 = bc[idir];
1915 
1916  int size = 2 * NC * ND * CEIL_NWP(Nst);
1917  int size_b = 2 * NC * ND2 * CEIL_NWP(Nxyz);
1918 
1919 #pragma acc data present(v2[0:size], buf[0:size_b]), copyin(bc2, Nxyz)
1920 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1921  {
1922 
1923 #pragma acc loop gang worker vector
1924  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
1925  int it = 0;
1926  int ist = ixyz + Nxyz * it;
1927 
1928  real_t wt1[2], wt2[2];
1929 
1930  for(int ic = 0; ic < NC; ++ic){
1931  int icr = 2 * ic;
1932  int ici = 2 * ic + 1;
1933 
1934  wt1[0] = bc2 * buf[icr + 2 * NVC * ixyz];
1935  wt1[1] = bc2 * buf[ici + 2 * NVC * ixyz];
1936  wt2[0] = bc2 * buf[icr + NVC + 2 * NVC * ixyz];
1937  wt2[1] = bc2 * buf[ici + NVC + 2 * NVC * ixyz];
1938 
1939  v2[IDX2_SP_R(ic, 0, ist)] += wt1[0];
1940  v2[IDX2_SP_I(ic, 0, ist)] += wt1[1];
1941  v2[IDX2_SP_R(ic, 1, ist)] += wt2[0];
1942  v2[IDX2_SP_I(ic, 1, ist)] += wt2[1];
1943  v2[IDX2_SP_R(ic, 2, ist)] -= wt1[0];
1944  v2[IDX2_SP_I(ic, 2, ist)] -= wt1[1];
1945  v2[IDX2_SP_R(ic, 3, ist)] -= wt2[0];
1946  v2[IDX2_SP_I(ic, 3, ist)] -= wt2[1];
1947  }
1948 
1949  }
1950 
1951  }
1952 
1953 }
1954 
1955 //====================================================================
1956 void mult_wilson_tmb_chiral(real_t *RESTRICT v2, real_t *RESTRICT u,
1957  real_t *RESTRICT v1,
1958  int *Nsize, int *bc, int Nc)
1959 {
1960  int idir = 3;
1961 
1962  int Nxyz = Nsize[0] * Nsize[1] * Nsize[2];
1963  int Nt = Nsize[3];
1964  int Nst = Nxyz * Nt;
1965 
1966  int size = 2 * NC * ND * CEIL_NWP(Nst);
1967  int size_u = NDF * CEIL_NWP(Nst);
1968 
1969 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]), \
1970  copyin(bc[0:4], idir, Nxyz, Nt, Nst)
1971 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
1972  {
1973 
1974 #pragma acc loop gang worker vector
1975  for(int ist = 0; ist < Nst; ++ist){
1976  int ixyz = ist % Nxyz;
1977  int it = ist/Nxyz;
1978  int it2 = (it-1+Nt) % Nt;
1979  int nei = ixyz + Nxyz * it2;
1980  real_t bc2 = 1.0;
1981  if(it == 0) bc2 = bc[idir];
1982 
1983  real_t vt[NVCD];
1984  real_t vt1[NVC], vt2[NVC], ut[NDF];
1985  real_t wt1[2], wt2[2];
1986 
1987  for(int id = 0; id < ND; ++id){
1988  for(int ic = 0; ic < NC; ++ic){
1989  int icr = 2*ic;
1990  int ici = 2*ic + 1;
1991  vt[icr + NVC * id] = v1[IDX2_SP_R(ic, id, nei)];
1992  vt[ici + NVC * id] = v1[IDX2_SP_I(ic, id, nei)];
1993  }
1994  }
1995 
1996  for(int ivc = 0; ivc < NVC; ++ivc){
1997  vt1[ivc] = vt[ivc + ID1] - vt[ivc + ID3];
1998  vt2[ivc] = vt[ivc + ID2] - vt[ivc + ID4];
1999  }
2000 
2001  for(int ic = 0; ic < NC; ++ic){
2002 
2003  for(int ic2 = 0; ic2 < NC; ++ic2){
2004  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, nei)];
2005  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, nei)];
2006  }
2007 
2008  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
2009  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
2010  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
2011  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
2012  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
2013  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
2014  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
2015  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
2016 
2017  v2[IDX2_SP_R(ic, 0, ist)] += bc2 * wt1[0];
2018  v2[IDX2_SP_I(ic, 0, ist)] += bc2 * wt1[1];
2019  v2[IDX2_SP_R(ic, 1, ist)] += bc2 * wt2[0];
2020  v2[IDX2_SP_I(ic, 1, ist)] += bc2 * wt2[1];
2021  v2[IDX2_SP_R(ic, 2, ist)] -= bc2 * wt1[0];
2022  v2[IDX2_SP_I(ic, 2, ist)] -= bc2 * wt1[1];
2023  v2[IDX2_SP_R(ic, 3, ist)] -= bc2 * wt2[0];
2024  v2[IDX2_SP_I(ic, 3, ist)] -= bc2 * wt2[1];
2025  }
2026 
2027  }
2028 
2029  }
2030 
2031 }
2032 
2033 //============================================================END=====
mult_wilson_xpb
void mult_wilson_xpb(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:130
mult_wilson_tpb_dirac
void mult_wilson_tpb_dirac(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1373
mult_wilson_yp1
void mult_wilson_yp1(real_t *RESTRICT buf, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:414
mult_wilson_tp2_dirac
void mult_wilson_tp2_dirac(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1311
mult_wilson_ypb
void mult_wilson_ypb(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:537
iy
int iy
Definition: mult_Wilson_xyz_openacc-inc.h:239
ixy
int ixy
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:513
mult_wilson_xmb
void mult_wilson_xmb(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:334
ID1
#define ID1
Definition: fopr_Wilson_impl_SU2-inc.h:18
izt
int izt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:262
mult_wilson_tmb_chiral
void mult_wilson_tmb_chiral(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1956
NVCD
#define NVCD
Definition: define_params_SU3.h:20
mult_wilson_zp1
void mult_wilson_zp1(real_t *RESTRICT buf, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:836
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
mult_wilson_xp1
void mult_wilson_xp1(real_t *RESTRICT buf, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:14
mult_wilson_ym2
void mult_wilson_ym2(real_t *RESTRICT v2, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:696
mult_wilson_ym1
void mult_wilson_ym1(real_t *RESTRICT buf, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:621
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
mult_wilson_tpb_chiral
void mult_wilson_tpb_chiral(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1759
IDX2_G_I
#define IDX2_G_I(ic1, ic2, ist)
Definition: define_index.h:52
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
mult_wilson_zm1
void mult_wilson_zm1(real_t *RESTRICT buf, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1045
idir
idir
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:264
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
mult_wilson_ymb
void mult_wilson_ymb(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:752
mult_wilson_zpb
void mult_wilson_zpb(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:962
mult_wilson_tp1_chiral
void mult_wilson_tp1_chiral(real_t *RESTRICT buf, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1647
ND
#define ND
Definition: field_F_imp_SU2-inc.h:18
mult_wilson_tm2_chiral
void mult_wilson_tm2_chiral(real_t *RESTRICT v2, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1905
it
int it
Definition: mult_Wilson_xyz_openacc-inc.h:461
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
mult_wilson_tm1_dirac
void mult_wilson_tm1_dirac(real_t *RESTRICT buf, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1450
mult_wilson_xm1
void mult_wilson_xm1(real_t *RESTRICT buf, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:209
ID3
#define ID3
Definition: fopr_Wilson_impl_SU2-inc.h:20
MULT_UV_I
#define MULT_UV_I(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_Wilson_dir_openacc-inc.h:11
mult_wilson_tp1_dirac
void mult_wilson_tp1_dirac(real_t *RESTRICT buf, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1261
mult_wilson_zmb
void mult_wilson_zmb(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1177
mult_wilson_xm2
void mult_wilson_xm2(real_t *RESTRICT v2, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:282
mult_wilson_zp2
void mult_wilson_zp2(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:891
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
mult_wilson_tmb_dirac
void mult_wilson_tmb_dirac(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1569
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
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
vt2
real_t vt2
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
mult_wilson_zm2
void mult_wilson_zm2(real_t *RESTRICT v2, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1121
mult_wilson_yp2
void mult_wilson_yp2(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:467
ND2
#define ND2
Definition: define_params_SU3.h:18
mult_wilson_tp2_chiral
void mult_wilson_tp2_chiral(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1693
mult_wilson_tm2_dirac
void mult_wilson_tm2_dirac(real_t *RESTRICT v2, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1522
mult_wilson_xp2
void mult_wilson_xp2(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:64
IDX2_SP_I
#define IDX2_SP_I(ic, id, ist)
Definition: define_index.h:32
mult_wilson_tm1_chiral
void mult_wilson_tm1_chiral(real_t *RESTRICT buf, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1837
MULT_UV_R
#define MULT_UV_R(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_Wilson_dir_openacc-inc.h:10