Bridge++  Ver.2.1.3
mult_Wilson_openacc-inc.h
Go to the documentation of this file.
1 
9 #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)
10 #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)
11 
12 #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)
13 #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)
14 #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)
15 #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)
16 
17 #define EXT_IMG_R(v1r,v1i,v2r,v2i,w1r,w1i,w2r,w2i) (v1r*w2r - v1i*w2i - v2r*w1r + v2i*w1i)
18 #define EXT_IMG_I(v1r,v1i,v2r,v2i,w1r,w1i,w2r,w2i) (- v1r*w2i - v1i*w2r + v2r*w1i + v2i*w1r)
19 
20 
21 //====================================================================
22 void mult_wilson_gm5_dirac(real_t *RESTRICT v2, real_t *RESTRICT v1,
23  int *Nsize, int Nc)
24 {
25  int Nst = Nsize[0] * Nsize[1] * Nsize[2] * Nsize[3];
26  int Nst_pad = CEIL_NWP(Nst);
27 
28  int size = NVC * ND * Nst_pad;
29 
30 #pragma acc data present(v2[0:size], v1[0:size]) copyin(Nst)
31 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
32  {
33 
34 #pragma acc loop gang worker vector
35  for(int ist = 0; ist < Nst; ++ist){
36  for(int ivc = 0; ivc < NVC; ++ivc){
37  v2[IDX2_SP(ivc, 0, ist)] = v1[IDX2_SP(ivc, 2, ist)];
38  v2[IDX2_SP(ivc, 1, ist)] = v1[IDX2_SP(ivc, 3, ist)];
39  v2[IDX2_SP(ivc, 2, ist)] = v1[IDX2_SP(ivc, 0, ist)];
40  v2[IDX2_SP(ivc, 3, ist)] = v1[IDX2_SP(ivc, 1, ist)];
41  }
42  }
43 
44  }
45 
46 }
47 
48 //====================================================================
49 void mult_wilson_gm5_chiral(real_t *RESTRICT v2, real_t *RESTRICT v1,
50  int *Nsize, int Nc)
51 {
52  int Nst = Nsize[0] * Nsize[1] * Nsize[2] * Nsize[3];
53  int Nst_pad = CEIL_NWP(Nst);
54 
55  int size = NVC * ND * Nst_pad;
56 
57 #pragma acc data present(v2[0:size], v1[0:size]) copyin(Nst)
58 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
59  {
60 
61 #pragma acc loop gang worker vector
62  for(int ist = 0; ist < Nst; ++ist){
63  for(int ivc = 0; ivc < NVC; ++ivc){
64  v2[IDX2_SP(ivc, 0, ist)] = v1[IDX2_SP(ivc, 0, ist)];
65  v2[IDX2_SP(ivc, 1, ist)] = v1[IDX2_SP(ivc, 1, ist)];
66  v2[IDX2_SP(ivc, 2, ist)] = -v1[IDX2_SP(ivc, 2, ist)];
67  v2[IDX2_SP(ivc, 3, ist)] = -v1[IDX2_SP(ivc, 3, ist)];
68  }
69  }
70 
71  }
72 
73 }
74 
75 //====================================================================
76 void mult_wilson_D_dirac(real_t *RESTRICT v2, real_t *RESTRICT u,
77  real_t *RESTRICT v1,
78  real_t kappa, int *Nsize, int *bc, int flag)
79 { // flag=0: D, flag=1: H
80 
81  int Nx = Nsize[0];
82  int Ny = Nsize[1];
83  int Nz = Nsize[2];
84  int Nt = Nsize[3];
85  int Nst = Nx * Ny * Nz * Nt;
86  int Nst_pad = CEIL_NWP(Nst);
87 
88  int nvst = NVC * ND * Nst_pad;
89  int ngst = NDF * Nst_pad * NDIM;
90 
91  real_t *RESTRICT u_up = u;
92  real_t *RESTRICT u_dn = u;
93 
94 #pragma acc data present(v1[0:nvst], u_up[0:ngst], u_dn[0:ngst], v2[0:nvst]) \
95  copyin(bc[0:4], kappa, Nx, Ny, Nz, Nt, Nst, Nst_pad)
96 
97 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
98  {
99 
100 #pragma acc loop gang worker vector
101  for(int site = 0; site < Nst; ++site){
102  int Nxy = Nx * Ny;
103  int Nxyz = Nx * Ny * Nz;
104 
105  real_t bc2;
106 
107  real_t u_0, u_1, u_2, u_3, u_4, u_5;
108  real_t u_6, u_7, u_8, u_9, u10, u11;
109  real_t u12, u13, u14, u15, u16, u17;
110  real_t vt1_0, vt1_1, vt1_2, vt1_3, vt1_4, vt1_5;
111  real_t vt2_0, vt2_1, vt2_2, vt2_3, vt2_4, vt2_5;
112  real_t wt1r, wt1i, wt2r, wt2i;
113 
114  real_t v2_01, v2_11, v2_21, v2_31, v2_41, v2_51;
115  real_t v2_02, v2_12, v2_22, v2_32, v2_42, v2_52;
116  real_t v2_03, v2_13, v2_23, v2_33, v2_43, v2_53;
117  real_t v2_04, v2_14, v2_24, v2_34, v2_44, v2_54;
118 
119 #include "inc/mult_Wilson_xyz_openacc-inc.h"
120 
121 #include "inc/mult_Wilson_t_dirac_openacc-inc.h"
122 
123  if(flag == 0){ // D
124  v2[IDX2_SP(0,0,site)] = v1[IDX2_SP(0,0,site)] - kappa * v2_01;
125  v2[IDX2_SP(1,0,site)] = v1[IDX2_SP(1,0,site)] - kappa * v2_11;
126  v2[IDX2_SP(2,0,site)] = v1[IDX2_SP(2,0,site)] - kappa * v2_21;
127  v2[IDX2_SP(3,0,site)] = v1[IDX2_SP(3,0,site)] - kappa * v2_31;
128  v2[IDX2_SP(4,0,site)] = v1[IDX2_SP(4,0,site)] - kappa * v2_41;
129  v2[IDX2_SP(5,0,site)] = v1[IDX2_SP(5,0,site)] - kappa * v2_51;
130 
131  v2[IDX2_SP(0,1,site)] = v1[IDX2_SP(0,1,site)] - kappa * v2_02;
132  v2[IDX2_SP(1,1,site)] = v1[IDX2_SP(1,1,site)] - kappa * v2_12;
133  v2[IDX2_SP(2,1,site)] = v1[IDX2_SP(2,1,site)] - kappa * v2_22;
134  v2[IDX2_SP(3,1,site)] = v1[IDX2_SP(3,1,site)] - kappa * v2_32;
135  v2[IDX2_SP(4,1,site)] = v1[IDX2_SP(4,1,site)] - kappa * v2_42;
136  v2[IDX2_SP(5,1,site)] = v1[IDX2_SP(5,1,site)] - kappa * v2_52;
137 
138  v2[IDX2_SP(0,2,site)] = v1[IDX2_SP(0,2,site)] - kappa * v2_03;
139  v2[IDX2_SP(1,2,site)] = v1[IDX2_SP(1,2,site)] - kappa * v2_13;
140  v2[IDX2_SP(2,2,site)] = v1[IDX2_SP(2,2,site)] - kappa * v2_23;
141  v2[IDX2_SP(3,2,site)] = v1[IDX2_SP(3,2,site)] - kappa * v2_33;
142  v2[IDX2_SP(4,2,site)] = v1[IDX2_SP(4,2,site)] - kappa * v2_43;
143  v2[IDX2_SP(5,2,site)] = v1[IDX2_SP(5,2,site)] - kappa * v2_53;
144 
145  v2[IDX2_SP(0,3,site)] = v1[IDX2_SP(0,3,site)] - kappa * v2_04;
146  v2[IDX2_SP(1,3,site)] = v1[IDX2_SP(1,3,site)] - kappa * v2_14;
147  v2[IDX2_SP(2,3,site)] = v1[IDX2_SP(2,3,site)] - kappa * v2_24;
148  v2[IDX2_SP(3,3,site)] = v1[IDX2_SP(3,3,site)] - kappa * v2_34;
149  v2[IDX2_SP(4,3,site)] = v1[IDX2_SP(4,3,site)] - kappa * v2_44;
150  v2[IDX2_SP(5,3,site)] = v1[IDX2_SP(5,3,site)] - kappa * v2_54;
151  }else{ // H
152  v2[IDX2_SP(0,2,site)] = v1[IDX2_SP(0,0,site)] - kappa * v2_01;
153  v2[IDX2_SP(1,2,site)] = v1[IDX2_SP(1,0,site)] - kappa * v2_11;
154  v2[IDX2_SP(2,2,site)] = v1[IDX2_SP(2,0,site)] - kappa * v2_21;
155  v2[IDX2_SP(3,2,site)] = v1[IDX2_SP(3,0,site)] - kappa * v2_31;
156  v2[IDX2_SP(4,2,site)] = v1[IDX2_SP(4,0,site)] - kappa * v2_41;
157  v2[IDX2_SP(5,2,site)] = v1[IDX2_SP(5,0,site)] - kappa * v2_51;
158 
159  v2[IDX2_SP(0,3,site)] = v1[IDX2_SP(0,1,site)] - kappa * v2_02;
160  v2[IDX2_SP(1,3,site)] = v1[IDX2_SP(1,1,site)] - kappa * v2_12;
161  v2[IDX2_SP(2,3,site)] = v1[IDX2_SP(2,1,site)] - kappa * v2_22;
162  v2[IDX2_SP(3,3,site)] = v1[IDX2_SP(3,1,site)] - kappa * v2_32;
163  v2[IDX2_SP(4,3,site)] = v1[IDX2_SP(4,1,site)] - kappa * v2_42;
164  v2[IDX2_SP(5,3,site)] = v1[IDX2_SP(5,1,site)] - kappa * v2_52;
165 
166  v2[IDX2_SP(0,0,site)] = v1[IDX2_SP(0,2,site)] - kappa * v2_03;
167  v2[IDX2_SP(1,0,site)] = v1[IDX2_SP(1,2,site)] - kappa * v2_13;
168  v2[IDX2_SP(2,0,site)] = v1[IDX2_SP(2,2,site)] - kappa * v2_23;
169  v2[IDX2_SP(3,0,site)] = v1[IDX2_SP(3,2,site)] - kappa * v2_33;
170  v2[IDX2_SP(4,0,site)] = v1[IDX2_SP(4,2,site)] - kappa * v2_43;
171  v2[IDX2_SP(5,0,site)] = v1[IDX2_SP(5,2,site)] - kappa * v2_53;
172 
173  v2[IDX2_SP(0,1,site)] = v1[IDX2_SP(0,3,site)] - kappa * v2_04;
174  v2[IDX2_SP(1,1,site)] = v1[IDX2_SP(1,3,site)] - kappa * v2_14;
175  v2[IDX2_SP(2,1,site)] = v1[IDX2_SP(2,3,site)] - kappa * v2_24;
176  v2[IDX2_SP(3,1,site)] = v1[IDX2_SP(3,3,site)] - kappa * v2_34;
177  v2[IDX2_SP(4,1,site)] = v1[IDX2_SP(4,3,site)] - kappa * v2_44;
178  v2[IDX2_SP(5,1,site)] = v1[IDX2_SP(5,3,site)] - kappa * v2_54;
179  }
180  }
181 
182  }
183 
184 }
185 
186 //====================================================================
187 void mult_wilson_D_chiral(real_t *RESTRICT v2, real_t *RESTRICT u,
188  real_t *RESTRICT v1,
189  real_t kappa, int *Nsize, int *bc, int flag)
190 { // flag=0: D, flag=1: H
191 
192  int Nx = Nsize[0];
193  int Ny = Nsize[1];
194  int Nz = Nsize[2];
195  int Nt = Nsize[3];
196  int Nst = Nx * Ny * Nz * Nt;
197  int Nst_pad = CEIL_NWP(Nst);
198 
199  int nvst = NVC * ND * Nst_pad;
200  int ngst = NDF * Nst_pad * NDIM;
201 
202  real_t *RESTRICT u_up = u;
203  real_t *RESTRICT u_dn = u;
204 
205 #pragma acc data present(v1[0:nvst], u_up[0:ngst], u_dn[0:ngst], v2[0:nvst]) \
206  copyin(bc[0:4], kappa, Nx, Ny, Nz, Nt,Nst, Nst_pad)
207 
208 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
209  {
210 
211 #pragma acc loop gang worker vector
212  for(int site = 0; site < Nst; ++site){
213  int Nxy = Nx * Ny;
214  int Nxyz = Nx * Ny * Nz;
215 
216  real_t bc2;
217 
218  real_t u_0, u_1, u_2, u_3, u_4, u_5;
219  real_t u_6, u_7, u_8, u_9, u10, u11;
220  real_t u12, u13, u14, u15, u16, u17;
221  real_t vt1_0, vt1_1, vt1_2, vt1_3, vt1_4, vt1_5;
222  real_t vt2_0, vt2_1, vt2_2, vt2_3, vt2_4, vt2_5;
223  real_t wt1r, wt1i, wt2r, wt2i;
224 
225  real_t v2_01, v2_11, v2_21, v2_31, v2_41, v2_51;
226  real_t v2_02, v2_12, v2_22, v2_32, v2_42, v2_52;
227  real_t v2_03, v2_13, v2_23, v2_33, v2_43, v2_53;
228  real_t v2_04, v2_14, v2_24, v2_34, v2_44, v2_54;
229 
230 #include "inc/mult_Wilson_xyz_openacc-inc.h"
231 
232 #include "inc/mult_Wilson_t_chiral_openacc-inc.h"
233 
234  if(flag == 0){ // D
235  // aypx and write back to global memory
236  v2[IDX2_SP(0,0,site)] = v1[IDX2_SP(0,0,site)] - kappa * v2_01;
237  v2[IDX2_SP(1,0,site)] = v1[IDX2_SP(1,0,site)] - kappa * v2_11;
238  v2[IDX2_SP(2,0,site)] = v1[IDX2_SP(2,0,site)] - kappa * v2_21;
239  v2[IDX2_SP(3,0,site)] = v1[IDX2_SP(3,0,site)] - kappa * v2_31;
240  v2[IDX2_SP(4,0,site)] = v1[IDX2_SP(4,0,site)] - kappa * v2_41;
241  v2[IDX2_SP(5,0,site)] = v1[IDX2_SP(5,0,site)] - kappa * v2_51;
242 
243  v2[IDX2_SP(0,1,site)] = v1[IDX2_SP(0,1,site)] - kappa * v2_02;
244  v2[IDX2_SP(1,1,site)] = v1[IDX2_SP(1,1,site)] - kappa * v2_12;
245  v2[IDX2_SP(2,1,site)] = v1[IDX2_SP(2,1,site)] - kappa * v2_22;
246  v2[IDX2_SP(3,1,site)] = v1[IDX2_SP(3,1,site)] - kappa * v2_32;
247  v2[IDX2_SP(4,1,site)] = v1[IDX2_SP(4,1,site)] - kappa * v2_42;
248  v2[IDX2_SP(5,1,site)] = v1[IDX2_SP(5,1,site)] - kappa * v2_52;
249 
250  v2[IDX2_SP(0,2,site)] = v1[IDX2_SP(0,2,site)] - kappa * v2_03;
251  v2[IDX2_SP(1,2,site)] = v1[IDX2_SP(1,2,site)] - kappa * v2_13;
252  v2[IDX2_SP(2,2,site)] = v1[IDX2_SP(2,2,site)] - kappa * v2_23;
253  v2[IDX2_SP(3,2,site)] = v1[IDX2_SP(3,2,site)] - kappa * v2_33;
254  v2[IDX2_SP(4,2,site)] = v1[IDX2_SP(4,2,site)] - kappa * v2_43;
255  v2[IDX2_SP(5,2,site)] = v1[IDX2_SP(5,2,site)] - kappa * v2_53;
256 
257  v2[IDX2_SP(0,3,site)] = v1[IDX2_SP(0,3,site)] - kappa * v2_04;
258  v2[IDX2_SP(1,3,site)] = v1[IDX2_SP(1,3,site)] - kappa * v2_14;
259  v2[IDX2_SP(2,3,site)] = v1[IDX2_SP(2,3,site)] - kappa * v2_24;
260  v2[IDX2_SP(3,3,site)] = v1[IDX2_SP(3,3,site)] - kappa * v2_34;
261  v2[IDX2_SP(4,3,site)] = v1[IDX2_SP(4,3,site)] - kappa * v2_44;
262  v2[IDX2_SP(5,3,site)] = v1[IDX2_SP(5,3,site)] - kappa * v2_54;
263  }else{ // H
264  v2[IDX2_SP(0,0,site)] = v1[IDX2_SP(0,0,site)] - kappa * v2_01;
265  v2[IDX2_SP(1,0,site)] = v1[IDX2_SP(1,0,site)] - kappa * v2_11;
266  v2[IDX2_SP(2,0,site)] = v1[IDX2_SP(2,0,site)] - kappa * v2_21;
267  v2[IDX2_SP(3,0,site)] = v1[IDX2_SP(3,0,site)] - kappa * v2_31;
268  v2[IDX2_SP(4,0,site)] = v1[IDX2_SP(4,0,site)] - kappa * v2_41;
269  v2[IDX2_SP(5,0,site)] = v1[IDX2_SP(5,0,site)] - kappa * v2_51;
270 
271  v2[IDX2_SP(0,1,site)] = v1[IDX2_SP(0,1,site)] - kappa * v2_02;
272  v2[IDX2_SP(1,1,site)] = v1[IDX2_SP(1,1,site)] - kappa * v2_12;
273  v2[IDX2_SP(2,1,site)] = v1[IDX2_SP(2,1,site)] - kappa * v2_22;
274  v2[IDX2_SP(3,1,site)] = v1[IDX2_SP(3,1,site)] - kappa * v2_32;
275  v2[IDX2_SP(4,1,site)] = v1[IDX2_SP(4,1,site)] - kappa * v2_42;
276  v2[IDX2_SP(5,1,site)] = v1[IDX2_SP(5,1,site)] - kappa * v2_52;
277 
278  v2[IDX2_SP(0,2,site)] = -v1[IDX2_SP(0,2,site)] + kappa * v2_03;
279  v2[IDX2_SP(1,2,site)] = -v1[IDX2_SP(1,2,site)] + kappa * v2_13;
280  v2[IDX2_SP(2,2,site)] = -v1[IDX2_SP(2,2,site)] + kappa * v2_23;
281  v2[IDX2_SP(3,2,site)] = -v1[IDX2_SP(3,2,site)] + kappa * v2_33;
282  v2[IDX2_SP(4,2,site)] = -v1[IDX2_SP(4,2,site)] + kappa * v2_43;
283  v2[IDX2_SP(5,2,site)] = -v1[IDX2_SP(5,2,site)] + kappa * v2_53;
284 
285  v2[IDX2_SP(0,3,site)] = -v1[IDX2_SP(0,3,site)] + kappa * v2_04;
286  v2[IDX2_SP(1,3,site)] = -v1[IDX2_SP(1,3,site)] + kappa * v2_14;
287  v2[IDX2_SP(2,3,site)] = -v1[IDX2_SP(2,3,site)] + kappa * v2_24;
288  v2[IDX2_SP(3,3,site)] = -v1[IDX2_SP(3,3,site)] + kappa * v2_34;
289  v2[IDX2_SP(4,3,site)] = -v1[IDX2_SP(4,3,site)] + kappa * v2_44;
290  v2[IDX2_SP(5,3,site)] = -v1[IDX2_SP(5,3,site)] + kappa * v2_54;
291  }
292  }
293 
294  }
295 
296 }
297 
298 //====================================================================
299 void mult_wilson_1_dirac(
300  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
301  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
302  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
303  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
304  real_t *RESTRICT u, real_t *RESTRICT v1,
305  int *Nsize, int *bc, int *do_comm)
306 {
307  int Nx = Nsize[0];
308  int Ny = Nsize[1];
309  int Nz = Nsize[2];
310  int Nt = Nsize[3];
311  int Nst = Nx * Ny * Nz * Nt;
312  int Nst_pad = CEIL_NWP(Nst);
313 
314  int size = NVC * ND * Nst_pad;
315  int size_u = NDF * Nst_pad * NDIM;
316  int size_bx = NVC * 2 * CEIL_NWP(Ny * Nz * Nt);
317  int size_by = NVC * 2 * CEIL_NWP(Nx * Nz * Nt);
318  int size_bz = NVC * 2 * CEIL_NWP(Nx * Ny * Nt);
319  int size_bt = NVC * 2 * CEIL_NWP(Nx * Ny * Nz);
320 
321 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
322  buf_yp[0:size_by], buf_ym[0:size_by], \
323  buf_zp[0:size_bz], buf_zm[0:size_bz], \
324  buf_tp[0:size_bt], buf_tm[0:size_bt], \
325  u[0:size_u], v1[0:size]) \
326  copyin(bc[0:4], do_comm[0:4], Nx, Ny, Nz, Nt, Nst_pad)
327 {
328 
329 #include "inc/mult_Wilson_1xyz_openacc-inc.h"
330 
331 
332  // idir = 3;
333  if(do_comm[3] > 0){
334 
335 #pragma acc parallel async \
336  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
337  {
338  int Nxyz = Nx * Ny * Nz;
339 
340 #pragma acc loop gang worker vector
341  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
342  int it = 0;
343  int ist = ixyz + Nxyz * it;
344  real_t bc2 = 2.0 * bc[3];
345  for(int ic = 0; ic < NC; ++ic){
346  buf_tp[IDXBF_R(ic, 0, ixyz)] = bc2 * v1[IDX2_SP_R(ic, 2, ist)];
347  buf_tp[IDXBF_I(ic, 0, ixyz)] = bc2 * v1[IDX2_SP_I(ic, 2, ist)];
348  buf_tp[IDXBF_R(ic, 1, ixyz)] = bc2 * v1[IDX2_SP_R(ic, 3, ist)];
349  buf_tp[IDXBF_I(ic, 1, ixyz)] = bc2 * v1[IDX2_SP_I(ic, 3, ist)];
350  }
351  }
352 
353  }
354 
355 #pragma acc parallel async \
356  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
357  {
358  int Nxyz = Nx * Ny * Nz;
359 
360 #pragma acc loop gang worker vector
361  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
362  int it = Nt-1;
363  int ist = ixyz + Nxyz * it;
364  int istu = ist + Nst_pad * 3;
365 
366  real_t vt1[NVC], vt2[NVC];
367  for(int ic = 0; ic < NC; ++ic){
368  int icr = 2*ic;
369  int ici = 2*ic + 1;
370  vt1[icr] = 2.0 * v1[IDX2_SP_R(ic, 0, ist)];
371  vt1[ici] = 2.0 * v1[IDX2_SP_I(ic, 0, ist)];
372  vt2[icr] = 2.0 * v1[IDX2_SP_R(ic, 1, ist)];
373  vt2[ici] = 2.0 * v1[IDX2_SP_I(ic, 1, ist)];
374  }
375 
376  for(int ic = 0; ic < NC; ++ic){
377 
378  real_t ut[NVC];
379  for(int ic2 = 0; ic2 < NC; ++ic2){
380  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
381  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
382  }
383 
384  real_t wt1[2], wt2[2];
385  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
386  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
387  wt1[1] = MULT_UV_I(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  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
390  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
391  wt2[1] = MULT_UV_I(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  buf_tm[IDXBF_R(ic, 0, ixyz)] = wt1[0];
394  buf_tm[IDXBF_I(ic, 0, ixyz)] = wt1[1];
395  buf_tm[IDXBF_R(ic, 1, ixyz)] = wt2[0];
396  buf_tm[IDXBF_I(ic, 1, ixyz)] = wt2[1];
397  }
398 
399  }
400 
401  }
402 
403  } // do_comm[3]
404 
405  #pragma acc wait
406 
407  } // acc data
408 
409  if(do_comm[0] > 0){
410  #pragma acc update async host (buf_xp[0:size_bx])
411  #pragma acc update async host (buf_xm[0:size_bx])
412  }
413  if(do_comm[1] > 0){
414  #pragma acc update async host (buf_yp[0:size_by])
415  #pragma acc update async host (buf_ym[0:size_by])
416  }
417  if(do_comm[2] > 0){
418  #pragma acc update async host (buf_zp[0:size_bz])
419  #pragma acc update async host (buf_zm[0:size_bz])
420  }
421  if(do_comm[3] > 0){
422  #pragma acc update async host (buf_tp[0:size_bt])
423  #pragma acc update async host (buf_tm[0:size_bt])
424  }
425 
426  #pragma acc wait
427 
428 }
429 
430 //====================================================================
431 void mult_wilson_1_chiral(
432  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
433  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
434  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
435  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
436  real_t *RESTRICT u, real_t *RESTRICT v1,
437  int *Nsize, int *bc, int *do_comm)
438 {
439  int Nx = Nsize[0];
440  int Ny = Nsize[1];
441  int Nz = Nsize[2];
442  int Nt = Nsize[3];
443  int Nst = Nx * Ny * Nz * Nt;
444  int Nst_pad = CEIL_NWP(Nst);
445 
446  int size = NVC * ND * Nst_pad;
447  int size_u = NDF * Nst_pad * NDIM;
448  int size_bx = NVC * 2 * CEIL_NWP(Ny * Nz * Nt);
449  int size_by = NVC * 2 * CEIL_NWP(Nx * Nz * Nt);
450  int size_bz = NVC * 2 * CEIL_NWP(Nx * Ny * Nt);
451  int size_bt = NVC * 2 * CEIL_NWP(Nx * Ny * Nz);
452 
453 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
454  buf_yp[0:size_by], buf_ym[0:size_by], \
455  buf_zp[0:size_bz], buf_zm[0:size_bz], \
456  buf_tp[0:size_bt], buf_tm[0:size_bt], \
457  u[0:size_u], v1[0:size]) \
458  copyin(bc[0:4], do_comm[0:4], Nx, Ny, Nz, Nt, Nst_pad)
459 {
460 
461 #include "inc/mult_Wilson_1xyz_openacc-inc.h"
462 
463 
464  // idir = 3;
465  if(do_comm[3] > 0){
466 
467 #pragma acc parallel async\
468  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
469  {
470  int Nxyz = Nx * Ny * Nz;
471 
472 #pragma acc loop gang worker vector
473  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
474  int it = 0;
475  int ist = ixyz + Nxyz * it;
476  real_t bc2 = bc[3];
477  for(int ivc = 0; ivc < NVC; ++ivc){
478  buf_tp[IDXBF(ivc, 0, ixyz)]
479  = bc2 * (v1[IDX2_SP(ivc, 0, ist)] + v1[IDX2_SP(ivc, 2, ist)]);
480  buf_tp[IDXBF(ivc, 1, ixyz)]
481  = bc2 * (v1[IDX2_SP(ivc, 1, ist)] + v1[IDX2_SP(ivc, 3, ist)]);
482  }
483  }
484 
485  }
486 
487 #pragma acc parallel async \
488  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
489  {
490  int Nxyz = Nx * Ny * Nz;
491 
492 #pragma acc loop gang worker vector
493  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
494  int it = Nt-1;
495  int ist = ixyz + Nxyz * it;
496  int istu = ist + Nst_pad * 3;
497 
498  real_t vt1[NVC], vt2[NVC];
499  for(int ivc = 0; ivc < NVC; ++ivc){
500  vt1[ivc] = v1[IDX2_SP(ivc, 0, ist)] - v1[IDX2_SP(ivc, 2, ist)];
501  vt2[ivc] = v1[IDX2_SP(ivc, 1, ist)] - v1[IDX2_SP(ivc, 3, ist)];
502  }
503 
504  for(int ic = 0; ic < NC; ++ic){
505 
506  real_t ut[NVC];
507  for(int ic2 = 0; ic2 < NC; ++ic2){
508  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
509  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
510  }
511 
512  real_t wt1[2], wt2[2];
513  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
514  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
515  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
516  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
517  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
518  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
519  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
520  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
521  buf_tm[IDXBF_R(ic, 0, ixyz)] = wt1[0];
522  buf_tm[IDXBF_I(ic, 0, ixyz)] = wt1[1];
523  buf_tm[IDXBF_R(ic, 1, ixyz)] = wt2[0];
524  buf_tm[IDXBF_I(ic, 1, ixyz)] = wt2[1];
525 
526  }
527 
528  }
529 
530  }
531 
532  } // do_comm[3]
533 
534  #pragma acc wait
535 
536  } // acc data
537 
538  if(do_comm[0] > 0){
539  #pragma acc update async host (buf_xp[0:size_bx])
540  #pragma acc update async host (buf_xm[0:size_bx])
541  }
542  if(do_comm[1] > 0){
543  #pragma acc update async host (buf_yp[0:size_by])
544  #pragma acc update async host (buf_ym[0:size_by])
545  }
546  if(do_comm[2] > 0){
547  #pragma acc update async host (buf_zp[0:size_bz])
548  #pragma acc update async host (buf_zm[0:size_bz])
549  }
550  if(do_comm[3] > 0){
551  #pragma acc update async host (buf_tp[0:size_bt])
552  #pragma acc update async host (buf_tm[0:size_bt])
553  }
554 
555  #pragma acc wait
556 
557 }
558 
559 //====================================================================
560 void mult_wilson_2_dirac(real_t *RESTRICT v2, real_t *RESTRICT u,
561  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
562  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
563  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
564  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
565  real_t kappa,
566  int *Nsize, int *bc, int *do_comm, int flag)
567 { // flag=0: D, flag=1: H
568 
569  int Nx = Nsize[0];
570  int Ny = Nsize[1];
571  int Nz = Nsize[2];
572  int Nt = Nsize[3];
573  int Nst = Nx * Ny * Nz * Nt;
574  int Nst_pad = CEIL_NWP(Nst);
575 
576  int size = NVC * ND * Nst_pad;
577  int size_u = NDF * Nst_pad * NDIM;
578  int size_bx = NVC * 2 * CEIL_NWP(Ny * Nz * Nt);
579  int size_by = NVC * 2 * CEIL_NWP(Nx * Nz * Nt);
580  int size_bz = NVC * 2 * CEIL_NWP(Nx * Ny * Nt);
581  int size_bt = NVC * 2 * CEIL_NWP(Nx * Ny * Nz);
582 
583  if(do_comm[0] > 0){
584  #pragma acc update async device (buf_xp[0:size_bx])
585  #pragma acc update async device (buf_xm[0:size_bx])
586  }
587  if(do_comm[1] > 0){
588  #pragma acc update async device (buf_yp[0:size_by])
589  #pragma acc update async device (buf_ym[0:size_by])
590  }
591  if(do_comm[2] > 0){
592  #pragma acc update async device (buf_zp[0:size_bz])
593  #pragma acc update async device (buf_zm[0:size_bz])
594  }
595  if(do_comm[3] > 0){
596  #pragma acc update async device (buf_tp[0:size_bt])
597  #pragma acc update async device (buf_tm[0:size_bt])
598  }
599 
600  #pragma acc wait
601 
602 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
603  buf_yp[0:size_by], buf_ym[0:size_by], \
604  buf_zp[0:size_bz], buf_zm[0:size_bz], \
605  buf_tp[0:size_bt], buf_tm[0:size_bt], \
606  v2[0:size], u[0:size_u]) \
607  copyin(bc[0:4], do_comm[0:4], kappa, Nx, Ny, Nz, Nt, Nst_pad)
608  {
609 
610 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
611  {
612  int Nst = Nx * Ny * Nz * Nt;
613 
614 #pragma acc loop gang worker vector
615  for(int ist = 0; ist < Nst; ++ist){
616  int ix = ist % Nx;
617  int iy = (ist/Nx) % Ny;
618  int iz = (ist/(Nx * Ny)) % Nz;
619  int it = ist/(Nx * Ny * Nz);
620 
621  real_t v2L[NVC*ND];
622  for(int id = 0; id < ND; ++id){
623  for(int ivc = 0; ivc < NVC; ++ivc){
624  v2L[ivc + NVC * id] = 0.0;
625  }
626  }
627 
628  int opr_any = 0;
629 
630 #include "inc/mult_Wilson_2xyz_openacc-inc.h"
631 
632  // idir = 3
633  if(do_comm[3] > 0){
634 
635  if(it == Nt-1){
636  int ixyz = ix + Nx * (iy + Ny * iz);
637  int istu = ist + Nst_pad * 3;
638 
639  real_t vt1[NVC], vt2[NVC];
640 
641  for(int ivc = 0; ivc < NVC; ++ivc){
642  vt1[ivc] = buf_tp[IDXBF(ivc, 0, ixyz)];
643  vt2[ivc] = buf_tp[IDXBF(ivc, 1, ixyz)];
644  }
645 
646  for(int ic = 0; ic < NC; ++ic){
647 
648  real_t ut[NVC];
649  for(int ic2 = 0; ic2 < NC; ++ic2){
650  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, istu)];
651  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, istu)];
652  }
653 
654  real_t wt1[2], wt2[2];
655  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
656  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
657  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
658  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
659  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
660  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
661  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
662  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
663  v2L[ic*2 + ID3] += wt1[0];
664  v2L[ic*2+1 + ID3] += wt1[1];
665  v2L[ic*2 + ID4] += wt2[0];
666  v2L[ic*2+1 + ID4] += wt2[1];
667  }
668  ++opr_any;
669  }
670 
671  if(it == 0){
672  int ixyz = ix + Nx * (iy + Ny * iz);
673  real_t bc2 = bc[3];
674  real_t wt1[2], wt2[2];
675 
676  for(int ic = 0; ic < NC; ++ic){
677  wt1[0] = bc2 * buf_tm[IDXBF_R(ic, 0, ixyz)];
678  wt1[1] = bc2 * buf_tm[IDXBF_I(ic, 0, ixyz)];
679  wt2[0] = bc2 * buf_tm[IDXBF_R(ic, 1, ixyz)];
680  wt2[1] = bc2 * buf_tm[IDXBF_I(ic, 1, ixyz)];
681  v2L[ic*2 + ID1] += wt1[0];
682  v2L[ic*2+1 + ID1] += wt1[1];
683  v2L[ic*2 + ID2] += wt2[0];
684  v2L[ic*2+1 + ID2] += wt2[1];
685  }
686  ++opr_any;
687  }
688 
689  }
690 
691  //axpy
692  if(opr_any > 0){
693  if(flag == 0){ //D
694  for(int id = 0; id < ND; ++id){
695  for(int ivc = 0; ivc < NVC; ++ivc){
696  v2[IDX2_SP(ivc, id, ist)] += -kappa * v2L[ivc + NVC * id];
697  }
698  }
699  }else{ // H
700  for(int ic = 0; ic < NC; ++ic){
701  v2[IDX2_SP_R(ic, 0, ist)] += -kappa * v2L[ic*2 + ID3];
702  v2[IDX2_SP_I(ic, 0, ist)] += -kappa * v2L[ic*2+1 + ID3];
703  v2[IDX2_SP_R(ic, 1, ist)] += -kappa * v2L[ic*2 + ID4];
704  v2[IDX2_SP_I(ic, 1, ist)] += -kappa * v2L[ic*2+1 + ID4];
705  v2[IDX2_SP_R(ic, 2, ist)] += -kappa * v2L[ic*2 + ID1];
706  v2[IDX2_SP_I(ic, 2, ist)] += -kappa * v2L[ic*2+1 + ID1];
707  v2[IDX2_SP_R(ic, 3, ist)] += -kappa * v2L[ic*2 + ID2];
708  v2[IDX2_SP_I(ic, 3, ist)] += -kappa * v2L[ic*2+1 + ID2];
709  }
710  }
711  }
712 
713  }
714 
715  }
716 
717  }
718 
719 }
720 
721 //====================================================================
722 void mult_wilson_2_chiral(real_t *RESTRICT v2, real_t *RESTRICT u,
723  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
724  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
725  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
726  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
727  real_t kappa,
728  int *Nsize, int *bc, int *do_comm, int flag)
729 {
730  int Nx = Nsize[0];
731  int Ny = Nsize[1];
732  int Nz = Nsize[2];
733  int Nt = Nsize[3];
734  int Nst = Nx * Ny * Nz * Nt;
735  int Nst_pad = CEIL_NWP(Nst);
736 
737  int size = NVC * ND * Nst_pad;
738  int size_u = NDF * Nst_pad * NDIM;
739  int size_bx = NVC * 2 * CEIL_NWP(Ny * Nz * Nt);
740  int size_by = NVC * 2 * CEIL_NWP(Nx * Nz * Nt);
741  int size_bz = NVC * 2 * CEIL_NWP(Nx * Ny * Nt);
742  int size_bt = NVC * 2 * CEIL_NWP(Nx * Ny * Nz);
743 
744  if(do_comm[0] > 0){
745  #pragma acc update async device (buf_xp[0:size_bx])
746  #pragma acc update async device (buf_xm[0:size_bx])
747  }
748  if(do_comm[1] > 0){
749  #pragma acc update async device (buf_yp[0:size_by])
750  #pragma acc update async device (buf_ym[0:size_by])
751  }
752  if(do_comm[2] > 0){
753  #pragma acc update async device (buf_zp[0:size_bz])
754  #pragma acc update async device (buf_zm[0:size_bz])
755  }
756  if(do_comm[3] > 0){
757  #pragma acc update async device (buf_tp[0:size_bt])
758  #pragma acc update async device (buf_tm[0:size_bt])
759  }
760 
761  #pragma acc wait
762 
763 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
764  buf_yp[0:size_by], buf_ym[0:size_by], \
765  buf_zp[0:size_bz], buf_zm[0:size_bz], \
766  buf_tp[0:size_bt], buf_tm[0:size_bt], \
767  v2[0:size], u[0:size_u]) \
768  copyin(bc[0:4], do_comm[0:4], kappa, Nx, Ny, Nz, Nt)
769  {
770 
771 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
772  {
773  int Nst = Nx * Ny * Nz * Nt;
774 
775 #pragma acc loop gang worker vector
776  for(int ist = 0; ist < Nst; ++ist){
777  int ix = ist % Nx;
778  int iy = (ist/Nx) % Ny;
779  int iz = (ist/(Nx * Ny)) % Nz;
780  int it = ist/(Nx * Ny * Nz);
781 
782  real_t v2L[NVC * ND];
783  for(int id = 0; id < ND; ++id){
784  for(int ivc = 0; ivc < NVC; ++ivc){
785  v2L[ivc + NVC * id] = 0.0;
786  }
787  }
788  int opr_any = 0;
789 
790 #include "inc/mult_Wilson_2xyz_openacc-inc.h"
791 
792  // idir = 3
793  if(do_comm[3] > 0){
794 
795  if(it == Nt-1){
796  int ixyz = ix + Nx * (iy + Ny * iz);
797  int istu = ist + Nst_pad * 3;
798 
799  real_t vt1[NVC], vt2[NVC], ut[NVC];
800  real_t wt1[2], wt2[2];
801 
802  for(int ivc = 0; ivc < NVC; ++ivc){
803  vt1[ivc] = buf_tp[IDXBF(ivc, 0, ixyz)];
804  vt2[ivc] = buf_tp[IDXBF(ivc, 1, ixyz)];
805  }
806 
807  for(int ic = 0; ic < NC; ++ic){
808 
809  for(int ic2 = 0; ic2 < NC; ++ic2){
810  ut[2*ic2 ] = u[IDX2_G_R(ic2, ic, istu)];
811  ut[2*ic2+1] = u[IDX2_G_I(ic2, ic, istu)];
812  }
813 
814  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
815  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
816  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
817  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
818  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
819  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
820  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
821  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
822 
823  v2L[ic*2 + ID1] += wt1[0];
824  v2L[ic*2+1 + ID1] += wt1[1];
825  v2L[ic*2 + ID2] += wt2[0];
826  v2L[ic*2+1 + ID2] += wt2[1];
827  v2L[ic*2 + ID3] += wt1[0];
828  v2L[ic*2+1 + ID3] += wt1[1];
829  v2L[ic*2 + ID4] += wt2[0];
830  v2L[ic*2+1 + ID4] += wt2[1];
831  }
832  ++opr_any;
833  }
834 
835  if(it == 0){
836  int ixyz = ix + Nx * (iy + Ny * iz);
837  real_t bc2 = bc[3];
838  real_t wt1[2], wt2[2];
839 
840  for(int ic = 0; ic < NC; ++ic){
841  wt1[0] = bc2 * buf_tm[IDXBF_R(ic, 0, ixyz)];
842  wt1[1] = bc2 * buf_tm[IDXBF_I(ic, 0, ixyz)];
843  wt2[0] = bc2 * buf_tm[IDXBF_R(ic, 1, ixyz)];
844  wt2[1] = bc2 * buf_tm[IDXBF_I(ic, 1, ixyz)];
845  v2L[ic*2 + ID1] += wt1[0];
846  v2L[ic*2+1 + ID1] += wt1[1];
847  v2L[ic*2 + ID2] += wt2[0];
848  v2L[ic*2+1 + ID2] += wt2[1];
849  v2L[ic*2 + ID3] -= wt1[0];
850  v2L[ic*2+1 + ID3] -= wt1[1];
851  v2L[ic*2 + ID4] -= wt2[0];
852  v2L[ic*2+1 + ID4] -= wt2[1];
853  }
854  ++opr_any;
855  }
856 
857  }
858 
859  if(opr_any > 0){
860  if(flag == 0){ // D
861  for(int id = 0; id < ND; ++id){
862  for(int ivc = 0; ivc < NVC; ++ivc){
863  v2[IDX2_SP(ivc, id, ist)] += -kappa * v2L[ivc + NVC * id];
864  }
865  }
866  }else{ // H
867  for(int ic = 0; ic < NC; ++ic){
868  v2[IDX2_SP_R(ic, 0, ist)] += -kappa * v2L[ic*2 + ID1];
869  v2[IDX2_SP_I(ic, 0, ist)] += -kappa * v2L[ic*2+1 + ID1];
870  v2[IDX2_SP_R(ic, 1, ist)] += -kappa * v2L[ic*2 + ID2];
871  v2[IDX2_SP_I(ic, 1, ist)] += -kappa * v2L[ic*2+1 + ID2];
872  v2[IDX2_SP_R(ic, 2, ist)] += kappa * v2L[ic*2 + ID3];
873  v2[IDX2_SP_I(ic, 2, ist)] += kappa * v2L[ic*2+1 + ID3];
874  v2[IDX2_SP_R(ic, 3, ist)] += kappa * v2L[ic*2 + ID4];
875  v2[IDX2_SP_I(ic, 3, ist)] += kappa * v2L[ic*2+1 + ID4];
876  }
877  }
878  }
879  }
880 
881  }
882 
883  }
884 
885 }
886 
887 //============================================================END=====
mult_wilson_1_dirac
void mult_wilson_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, int *Nsize, int *bc, int *do_comm)
Definition: mult_Wilson_openacc-inc.h:299
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
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
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
mult_wilson_2_dirac
void mult_wilson_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, int *Nsize, int *bc, int *do_comm, int flag)
Definition: mult_Wilson_openacc-inc.h:560
MULT_UV_I
#define MULT_UV_I(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_Wilson_openacc-inc.h:10
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
ID1
#define ID1
Definition: fopr_Wilson_impl_SU2-inc.h:18
mult_wilson_gm5_chiral
void mult_wilson_gm5_chiral(real_t *RESTRICT v2, real_t *RESTRICT v1, int *Nsize, int Nc)
Definition: mult_Wilson_openacc-inc.h:49
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
MULT_UV_R
#define MULT_UV_R(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_Wilson_openacc-inc.h:9
mult_wilson_D_dirac
void mult_wilson_D_dirac(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, real_t kappa, int *Nsize, int *bc, int flag)
Definition: mult_Wilson_openacc-inc.h:76
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
mult_Wilson_t_dirac_openacc-inc.h
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
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
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_2_chiral
void mult_wilson_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, int *Nsize, int *bc, int *do_comm, int flag)
Definition: mult_Wilson_openacc-inc.h:722
IDXBF
#define IDXBF(ivc, id, ist)
Definition: define_index.h:43
mult_Wilson_t_chiral_openacc-inc.h
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
mult_Wilson_2xyz_openacc-inc.h
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
v2_51
v2_51
Definition: mult_Clover_csw_chiral_openacc-inc.h:87
mult_wilson_1_chiral
void mult_wilson_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, int *Nsize, int *bc, int *do_comm)
Definition: mult_Wilson_openacc-inc.h:431
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
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
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
NDIM
#define NDIM
Definition: contract_4spinor.cpp:18
mult_wilson_D_chiral
void mult_wilson_D_chiral(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, real_t kappa, int *Nsize, int *bc, int flag)
Definition: mult_Wilson_openacc-inc.h:187
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
mult_wilson_gm5_dirac
void mult_wilson_gm5_dirac(real_t *RESTRICT v2, real_t *RESTRICT v1, int *Nsize, int Nc)
Definition: mult_Wilson_openacc-inc.h:22
vt2
real_t vt2
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
IDX2_SP_I
#define IDX2_SP_I(ic, id, ist)
Definition: define_index.h:32
mult_Wilson_xyz_openacc-inc.h
u_7
u_7
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:70
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
mult_Wilson_1xyz_openacc-inc.h
vt1_1
vt1_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:33