Bridge++  Ver.2.1.3
mult_Domainwall_5din_eo_openacc-inc.h
Go to the documentation of this file.
1 
10 #ifndef MULT_DOMAINWALL_5DIN_EO_ACC_INCLUDED
11 #define MULT_DOMAINWALL_5DIN_EO_ACC_INCLUDED
12 
13 //====================================================================
14 void mult_domainwall_5din_ee_5dir_dirac(
15  real_t *RESTRICT vp, real_t *RESTRICT wp,
16  real_t mq, real_t M0, int Ns, real_t *b, real_t *c, real_t alpha,
17  int *Nsize)
18 {
19  int Nx = Nsize[0];
20  int Ny = Nsize[1];
21  int Nz = Nsize[2];
22  int Nt = Nsize[3];
23  int Nst = Nx * Ny * Nz * Nt;
24  int Nst_pad = CEIL_NWP(Nst);
25 
26  int Nin5 = NVCD * Ns;
27  int size = Nin5 * Nst_pad;
28 
29 #pragma acc data present(vp[0:size], wp[0:size]) \
30  copyin(Nst, Nst_pad, Ns, Nin5, mq, M0, b[0:Ns], c[0:Ns], alpha)
31  {
32 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
33  {
34 
35 #pragma acc loop gang worker vector
36  for (int site = 0; site < Nst_pad; ++site) {
37  if(site < Nst){
38 
39  for (int is = 0; is < Ns; ++is) {
40 
41  real_t FF1 = b[is] * (4.0 - M0) + 1.0;
42  real_t FF2 = c[is] * (4.0 - M0) - 1.0;
43 
44  real_t vt[NVCD], wt[NVCD];
45 
46  int is_up = (is+1) % Ns;
47  real_t Fup = 0.5 * FF2 * alpha;
48  if (is == Ns-1) Fup = -0.5 * mq * FF2;
49 
50  for (int ivcd = 0; ivcd < NVCD; ++ivcd) {
51  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is_up), site)];
52  }
53  for (int ivc = 0; ivc < NVC; ++ivc) {
54  vt[ID1 + ivc] = Fup * (wt[ID1 + ivc] - wt[ID3 + ivc]);
55  vt[ID2 + ivc] = Fup * (wt[ID2 + ivc] - wt[ID4 + ivc]);
56  vt[ID3 + ivc] = Fup * (wt[ID3 + ivc] - wt[ID1 + ivc]);
57  vt[ID4 + ivc] = Fup * (wt[ID4 + ivc] - wt[ID2 + ivc]);
58  }
59 
60  int is_dn = (is-1 + Ns) % Ns;
61  real_t Fdn = 0.5 * FF2 * alpha;
62  if (is == 0) Fdn = -0.5 * mq * FF2;
63 
64  for (int ivcd = 0; ivcd < NVCD; ++ivcd) {
65  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is_dn), site)];
66  }
67  for (int ivc = 0; ivc < NVC; ++ivc) {
68  vt[ID1 + ivc] += Fdn * (wt[ID1 + ivc] + wt[ID3 + ivc]);
69  vt[ID2 + ivc] += Fdn * (wt[ID2 + ivc] + wt[ID4 + ivc]);
70  vt[ID3 + ivc] += Fdn * (wt[ID3 + ivc] + wt[ID1 + ivc]);
71  vt[ID4 + ivc] += Fdn * (wt[ID4 + ivc] + wt[ID2 + ivc]);
72  }
73 
74  if(is == 0){
75  real_t fac1 = FF1 * 0.5 * ( 1.0 + alpha);
76  real_t fac2 = FF1 * 0.5 * (-1.0 + alpha);
77  for(int ivc = 0; ivc < NVC; ++ivc){
78  real_t wt1, wt2, wt3, wt4;
79  wt1 = wp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)];
80  wt2 = wp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)];
81  wt3 = wp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)];
82  wt4 = wp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)];
83  wt[ID1 + ivc] = fac1 * wt1 + fac2 * wt3;
84  wt[ID2 + ivc] = fac1 * wt2 + fac2 * wt4;
85  wt[ID3 + ivc] = fac1 * wt3 + fac2 * wt1;
86  wt[ID4 + ivc] = fac1 * wt4 + fac2 * wt2;
87  }
88  }else if(is == Ns-1){
89  real_t fac1 = FF1 * 0.5 * (1.0 + alpha);
90  real_t fac2 = FF1 * 0.5 * (1.0 - alpha);
91  for(int ivc = 0; ivc < NVC; ++ivc){
92  real_t wt1, wt2, wt3, wt4;
93  wt1 = wp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)];
94  wt2 = wp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)];
95  wt3 = wp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)];
96  wt4 = wp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)];
97  wt[ID1 + ivc] = fac1 * wt1 + fac2 * wt3;
98  wt[ID2 + ivc] = fac1 * wt2 + fac2 * wt4;
99  wt[ID3 + ivc] = fac1 * wt3 + fac2 * wt1;
100  wt[ID4 + ivc] = fac1 * wt4 + fac2 * wt2;
101  }
102  }else{
103  real_t f1 = FF1 * alpha;
104  for (int ivcd = 0; ivcd < NVCD; ++ivcd) {
105  wt[ivcd] = f1 * wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
106  }
107  }
108 
109  for (int ivcd = 0; ivcd < NVCD; ++ivcd) {
110  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = wt[ivcd] + vt[ivcd];
111  }
112  }
113  }
114  } // site loop
115 
116  }
117  }
118 
119 }
120 
121 //====================================================================
122 void mult_domainwall_5din_eo_5dir_dirac(
123  real_t *RESTRICT yp, real_t *RESTRICT wp,
124  real_t mq, real_t M0, int Ns, real_t *b, real_t *c, real_t alpha,
125  int *Nsize)
126 {
127  int Nx = Nsize[0];
128  int Ny = Nsize[1];
129  int Nz = Nsize[2];
130  int Nt = Nsize[3];
131  int Nst = Nx * Ny * Nz * Nt;
132  int Nst_pad = CEIL_NWP(Nst);
133 
134  int Nin5 = NVCD * Ns;
135  int size = Nin5 * Nst_pad;
136 
137 #pragma acc data present(yp[0:size], wp[0:size]) \
138  copyin(Nst, Nst_pad, Ns, Nin5, mq, M0, b[0:Ns], c[0:Ns], alpha)
139  {
140 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
141  {
142 
143 #pragma acc loop gang worker vector
144  for (int idx = 0; idx < Nst_pad * NVC; ++idx) {
145  int idx2_wp = idx / NWP;
146  int idx_in = idx % NWP;
147  int ivc = idx2_wp % NVC;
148  int idx_out = idx2_wp / NVC;
149  int site = idx_in + NWP*idx_out;
150  if(site < Nst){
151 
152  for (int is = 0; is < Ns; ++is) {
153 
154  real_t vt1, vt2, vt3, vt4;
155  real_t wt1, wt2, wt3, wt4;
156 
157  int is_up = (is+1) % Ns;
158  real_t Fup = 0.5 * c[is] * alpha;
159  if (is == Ns-1) Fup = -0.5 * mq * c[is];
160  wt1 = wp[IDX2(Nin5, (ID1 + ivc + NVCD * is_up), site)];
161  wt2 = wp[IDX2(Nin5, (ID2 + ivc + NVCD * is_up), site)];
162  wt3 = wp[IDX2(Nin5, (ID3 + ivc + NVCD * is_up), site)];
163  wt4 = wp[IDX2(Nin5, (ID4 + ivc + NVCD * is_up), site)];
164 
165  vt1 = Fup * (wt1 - wt3);
166  vt2 = Fup * (wt2 - wt4);
167  vt3 = Fup * (wt3 - wt1);
168  vt4 = Fup * (wt4 - wt2);
169 
170  int is_dn = (is-1 + Ns) % Ns;
171  real_t Fdn = 0.5 * c[is] * alpha;
172  if (is == 0) Fdn = -0.5 * mq * c[is];
173  wt1 = wp[IDX2(Nin5, (ID1 + ivc + NVCD * is_dn), site)];
174  wt2 = wp[IDX2(Nin5, (ID2 + ivc + NVCD * is_dn), site)];
175  wt3 = wp[IDX2(Nin5, (ID3 + ivc + NVCD * is_dn), site)];
176  wt4 = wp[IDX2(Nin5, (ID4 + ivc + NVCD * is_dn), site)];
177 
178  vt1 += Fdn * (wt1+ wt3);
179  vt2 += Fdn * (wt2+ wt4);
180  vt3 += Fdn * (wt3+ wt1);
181  vt4 += Fdn * (wt4+ wt2);
182 
183  if(is == 0){
184  real_t b1 = b[is] * 0.5 * ( 1.0 + alpha);
185  real_t b2 = b[is] * 0.5 * (-1.0 + alpha);
186  real_t yt1, yt2, yt3, yt4;
187  yt1 = wp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)];
188  yt2 = wp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)];
189  yt3 = wp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)];
190  yt4 = wp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)];
191  wt1 = b1 * yt1 + b2 * yt3;
192  wt2 = b1 * yt2 + b2 * yt4;
193  wt3 = b1 * yt3 + b2 * yt1;
194  wt4 = b1 * yt4 + b2 * yt2;
195  }else if(is == Ns-1){
196  real_t b1 = b[is] * 0.5 * (1.0 + alpha);
197  real_t b2 = b[is] * 0.5 * (1.0 - alpha);
198  real_t yt1, yt2, yt3, yt4;
199  yt1 = wp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)];
200  yt2 = wp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)];
201  yt3 = wp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)];
202  yt4 = wp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)];
203  wt1 = b1 * yt1 + b2 * yt3;
204  wt2 = b1 * yt2 + b2 * yt4;
205  wt3 = b1 * yt3 + b2 * yt1;
206  wt4 = b1 * yt4 + b2 * yt2;
207  }else{
208  real_t bb = b[is] * alpha;
209  wt1 = bb * wp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)];
210  wt2 = bb * wp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)];
211  wt3 = bb * wp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)];
212  wt4 = bb * wp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)];
213  }
214 
215  yp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)] = -0.5 * (wt1 + vt1);
216  yp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)] = -0.5 * (wt2 + vt2);
217  yp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)] = -0.5 * (wt3 + vt3);
218  yp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)] = -0.5 * (wt4 + vt4);
219 
220  } //is
221 
222  } // site < Nst
223  } // idx
224  }
225 
226  }
227 }
228 
229 //====================================================================
230 void mult_domainwall_5din_ee_5dirdag_dirac(
231  real_t *RESTRICT vp, real_t *RESTRICT wp,
232  real_t mq, real_t M0, int Ns, real_t *b, real_t *c, real_t alpha,
233  int *Nsize)
234 {
235  int Nx = Nsize[0];
236  int Ny = Nsize[1];
237  int Nz = Nsize[2];
238  int Nt = Nsize[3];
239  int Nst = Nx * Ny * Nz * Nt;
240  int Nst_pad = CEIL_NWP(Nst);
241 
242  int Nin5 = NVCD * Ns;
243  int size = Nin5 * Nst_pad;
244 
245 #pragma acc data present(vp[0:size], wp[0:size]) \
246  copyin(Nst, Nst_pad, Ns, Nin5, mq, M0, b[0:Ns], c[0:Ns], alpha)
247  {
248 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
249  {
250 
251 #pragma acc loop gang worker vector
252  for (int site = 0; site < Nst_pad; ++site) {
253  if(site < Nst){
254 
255  for (int is = 0; is < Ns; ++is) {
256 
257  real_t vt[NVCD], xt[NVCD];
258 
259  real_t B1 = b[is] * (4.0 - M0) + 1.0;
260  real_t a1 = -0.5 * b[is];
261 
262  if(is == 0){
263  real_t fac1 = B1 * 0.5 * ( 1.0 + alpha);
264  real_t fac2 = B1 * 0.5 * (-1.0 + alpha);
265  for(int ivc = 0; ivc < NVC; ++ivc){
266  real_t wt1, wt2, wt3, wt4;
267  wt1 = wp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)];
268  wt2 = wp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)];
269  wt3 = wp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)];
270  wt4 = wp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)];
271  vt[ID1 + ivc] = fac1 * wt1 + fac2 * wt3;
272  vt[ID2 + ivc] = fac1 * wt2 + fac2 * wt4;
273  vt[ID3 + ivc] = fac1 * wt3 + fac2 * wt1;
274  vt[ID4 + ivc] = fac1 * wt4 + fac2 * wt2;
275  }
276  }else if(is == Ns-1){
277  real_t fac1 = B1 * 0.5 * (1.0 + alpha);
278  real_t fac2 = B1 * 0.5 * (1.0 - alpha);
279  for(int ivc = 0; ivc < NVC; ++ivc){
280  real_t wt1, wt2, wt3, wt4;
281  wt1 = wp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)];
282  wt2 = wp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)];
283  wt3 = wp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)];
284  wt4 = wp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)];
285  vt[ID1 + ivc] = fac1 * wt1 + fac2 * wt3;
286  vt[ID2 + ivc] = fac1 * wt2 + fac2 * wt4;
287  vt[ID3 + ivc] = fac1 * wt3 + fac2 * wt1;
288  vt[ID4 + ivc] = fac1 * wt4 + fac2 * wt2;
289  }
290  }else{
291  real_t bb = B1 * alpha;
292  for (int ivcd = 0; ivcd < NVCD; ++ivcd) {
293  vt[ivcd] = bb * wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
294  }
295  }
296 
297  int is_up = (is+1) % Ns;
298  for (int ivcd = 0; ivcd < NVCD; ++ivcd) {
299  xt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is_up), site)];
300  }
301 
302  real_t Fup = 0.5 * (c[is_up] * (4.0 - M0) - 1.0);
303  if (is == Ns-1){
304  Fup *= -mq;
305  }else{
306  Fup *= alpha;
307  }
308  for (int ivc = 0; ivc < NVC; ++ivc) {
309  vt[ID1 + ivc] += Fup * (xt[ID1 + ivc] + xt[ID3 + ivc]);
310  vt[ID2 + ivc] += Fup * (xt[ID2 + ivc] + xt[ID4 + ivc]);
311  vt[ID3 + ivc] += Fup * (xt[ID3 + ivc] + xt[ID1 + ivc]);
312  vt[ID4 + ivc] += Fup * (xt[ID4 + ivc] + xt[ID2 + ivc]);
313  }
314 
315  int is_dn = (is-1 + Ns) % Ns;
316  for (int ivcd = 0; ivcd < NVCD; ++ivcd) {
317  xt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is_dn), site)];
318  }
319 
320  real_t Fdn = 0.5 * (c[is_dn] * (4.0 - M0) - 1.0);
321  if (is == 0){
322  Fdn *= -mq;
323  }else{
324  Fdn *= alpha;
325  }
326  for (int ivc = 0; ivc < NVC; ++ivc) {
327  vt[ID1 + ivc] += Fdn * (xt[ID1 + ivc] - xt[ID3 + ivc]);
328  vt[ID2 + ivc] += Fdn * (xt[ID2 + ivc] - xt[ID4 + ivc]);
329  vt[ID3 + ivc] += Fdn * (xt[ID3 + ivc] - xt[ID1 + ivc]);
330  vt[ID4 + ivc] += Fdn * (xt[ID4 + ivc] - xt[ID2 + ivc]);
331  }
332 
333  for (int ivcd = 0; ivcd < NVCD; ++ivcd) {
334  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vt[ivcd];
335  }
336  }
337  }
338  }
339 
340  }
341  }
342 
343 }
344 
345 //====================================================================
346 void mult_domainwall_5din_eo_5dirdag_dirac(
347  real_t *RESTRICT vp, real_t *RESTRICT yp,
348  real_t mq, real_t M0, int Ns, real_t *b, real_t *c, real_t alpha,
349  int *Nsize)
350 {
351  int Nx = Nsize[0];
352  int Ny = Nsize[1];
353  int Nz = Nsize[2];
354  int Nt = Nsize[3];
355  int Nst = Nx * Ny * Nz * Nt;
356  int Nst_pad = CEIL_NWP(Nst);
357 
358  int Nin5 = NVCD * Ns;
359  int size = Nin5 * Nst_pad;
360 
361 #pragma acc data present(vp[0:size], yp[0:size]) \
362  copyin(Nst, Nst_pad, Ns, Nin5, mq, M0, b[0:Ns], c[0:Ns], alpha)
363  {
364 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
365  {
366 
367 #pragma acc loop gang worker vector
368  for (int idx = 0; idx < NVC * Nst_pad; ++idx) {
369  int idx2_wp = idx / NWP;
370  int idx_in = idx % NWP;
371  int ivc = idx2_wp % NVC;
372  int idx_out = idx2_wp / NVC;
373  int site = idx_in + NWP*idx_out;
374  if(site<Nst){
375 
376  for (int is = 0; is < Ns; ++is) {
377 
378  real_t vt1, vt2, vt3, vt4;
379  real_t yt1, yt2, yt3, yt4;
380 
381  yt1 = yp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)];
382  yt2 = yp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)];
383  yt3 = yp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)];
384  yt4 = yp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)];
385 
386  if(is == 0){
387  real_t b1 = -0.5 * b[is] * 0.5 * ( 1.0 + alpha);
388  real_t b2 = -0.5 * b[is] * 0.5 * (-1.0 + alpha);
389  vt1 = b1 * yt3 + b2 * yt1;
390  vt2 = b1 * yt4 + b2 * yt2;
391  vt3 = b1 * yt1 + b2 * yt3;
392  vt4 = b1 * yt2 + b2 * yt4;
393  }else if(is == Ns-1){
394  real_t b1 = -0.5 * b[is] * 0.5 * (1.0 + alpha);
395  real_t b2 = -0.5 * b[is] * 0.5 * (1.0 - alpha);
396  vt1 = b1 * yt3 + b2 * yt1;
397  vt2 = b1 * yt4 + b2 * yt2;
398  vt3 = b1 * yt1 + b2 * yt3;
399  vt4 = b1 * yt2 + b2 * yt4;
400  }else{
401  real_t bb = -0.5 * b[is] * alpha;
402  vt1 = bb * yt3;
403  vt2 = bb * yt4;
404  vt3 = bb * yt1;
405  vt4 = bb * yt2;
406  }
407 
408  int is_up = (is+1) % Ns;
409  yt1 = yp[IDX2(Nin5, (ID1 + ivc + NVCD * is_up), site)];
410  yt2 = yp[IDX2(Nin5, (ID2 + ivc + NVCD * is_up), site)];
411  yt3 = yp[IDX2(Nin5, (ID3 + ivc + NVCD * is_up), site)];
412  yt4 = yp[IDX2(Nin5, (ID4 + ivc + NVCD * is_up), site)];
413 
414  real_t Fup = -0.5 * c[is_up] * 0.5 * alpha;
415  if (is == Ns-1) Fup = -0.5 * c[is_up] * (-0.5) * mq;
416 
417  vt1 += Fup * (yt3 + yt1);
418  vt2 += Fup * (yt4 + yt2);
419  vt3 += Fup * (yt1 + yt3);
420  vt4 += Fup * (yt2 + yt4);
421 
422  int is_dn = (is-1 + Ns) % Ns;
423  yt1 = yp[IDX2(Nin5, (ID1 + ivc + NVCD * is_dn), site)];
424  yt2 = yp[IDX2(Nin5, (ID2 + ivc + NVCD * is_dn), site)];
425  yt3 = yp[IDX2(Nin5, (ID3 + ivc + NVCD * is_dn), site)];
426  yt4 = yp[IDX2(Nin5, (ID4 + ivc + NVCD * is_dn), site)];
427 
428  real_t Fdn = -0.5 * c[is_dn] * (0.5) * alpha;
429  if (is == 0) Fdn = -0.5 * c[is_dn] * (-0.5) * mq;
430 
431  vt1 += Fdn * (yt3 - yt1);
432  vt2 += Fdn * (yt4 - yt2);
433  vt3 += Fdn * (yt1 - yt3);
434  vt4 += Fdn * (yt2 - yt4);
435 
436  vp[IDX2(Nin5, (ID1 + ivc + NVCD * is), site)] = vt1;
437  vp[IDX2(Nin5, (ID2 + ivc + NVCD * is), site)] = vt2;
438  vp[IDX2(Nin5, (ID3 + ivc + NVCD * is), site)] = vt3;
439  vp[IDX2(Nin5, (ID4 + ivc + NVCD * is), site)] = vt4;
440  } // is
441 
442  } // site < Nst
443  } // idx
444  }
445 
446 }
447 }
448 
449 //====================================================================
450 void mult_domainwall_5din_eo_hopb_dirac_5d(
451  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
452  int Ns, int *bc, int *Nsize, int *do_comm, int ieo, int jeo,
453  int jgm5)
454 {
455  int Nx = Nsize[0];
456  int Ny = Nsize[1];
457  int Nz = Nsize[2];
458  int Nt = Nsize[3];
459  int Nst = Nx * Ny * Nz * Nt;
460  int Nst_pad = CEIL_NWP(Nst);
461 
462  int Nin5 = NVCD * Ns;
463 
464  int size = Nin5 * Nst_pad;
465  int size_u = NDF * Nst_pad * 2 * NDIM;
466 
467 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
468  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
469  bc[0:NDIM], do_comm[0:NDIM], ieo, jeo, jgm5)
470  {
471 
472 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
473  {
474  real_t *RESTRICT u_up = up;
475  real_t *RESTRICT u_dn = up;
476 
477 #pragma acc loop gang worker vector
478  for (int idx = 0; idx < Ns * Nst_pad; ++idx) {
479  int idx2_wp = idx / NWP;
480  int idx_in = idx % NWP;
481  int is = idx2_wp % Ns;
482  int idx_out = idx2_wp / Ns;
483 
484  int site = idx_in + NWP*idx_out;
485  if(site < Nst) {
486 
487  int Nxy = Nx * Ny;
488  int ix = site % Nx;
489  int iyzt = site / Nx;
490  int iy = iyzt % Ny;
491  int izt = site / Nxy;
492  int iz = izt % Nz;
493  int it = izt / Nz;
494  int Nxyz = Nx * Ny * Nz;
495 
496  int keo = (jeo + iy + iz + it) % 2;
497 
498  real_t u_0, u_1, u_2, u_3, u_4, u_5;
499  real_t u_6, u_7, u_8, u_9, u10, u11;
500  real_t u12, u13, u14, u15, u16, u17;
501  real_t vt1_0, vt1_1, vt1_2, vt1_3, vt1_4, vt1_5;
502  real_t vt2_0, vt2_1, vt2_2, vt2_3, vt2_4, vt2_5;
503  real_t wt1r, wt1i, wt2r, wt2i;
504 
505  real_t v2_01, v2_11, v2_21, v2_31, v2_41, v2_51;
506  real_t v2_02, v2_12, v2_22, v2_32, v2_42, v2_52;
507  real_t v2_03, v2_13, v2_23, v2_33, v2_43, v2_53;
508  real_t v2_04, v2_14, v2_24, v2_34, v2_44, v2_54;
509 
510  real_t *RESTRICT v1 = wp;
511  real_t *RESTRICT v2 = vp;
512 
513 #include "inc/mult_Domainwall_eo_xyz_openacc-inc.h"
514 
515 #include "inc/mult_Domainwall_eo_t_dirac_openacc-inc.h"
516 
517  v2[IDX2_SP_5D_R(0,0,is,Ns,site)] = v2_01;
518  v2[IDX2_SP_5D_I(0,0,is,Ns,site)] = v2_11;
519  v2[IDX2_SP_5D_R(1,0,is,Ns,site)] = v2_21;
520  v2[IDX2_SP_5D_I(1,0,is,Ns,site)] = v2_31;
521  v2[IDX2_SP_5D_R(2,0,is,Ns,site)] = v2_41;
522  v2[IDX2_SP_5D_I(2,0,is,Ns,site)] = v2_51;
523 
524  v2[IDX2_SP_5D_R(0,1,is,Ns,site)] = v2_02;
525  v2[IDX2_SP_5D_I(0,1,is,Ns,site)] = v2_12;
526  v2[IDX2_SP_5D_R(1,1,is,Ns,site)] = v2_22;
527  v2[IDX2_SP_5D_I(1,1,is,Ns,site)] = v2_32;
528  v2[IDX2_SP_5D_R(2,1,is,Ns,site)] = v2_42;
529  v2[IDX2_SP_5D_I(2,1,is,Ns,site)] = v2_52;
530 
531  v2[IDX2_SP_5D_R(0,2,is,Ns,site)] = v2_03;
532  v2[IDX2_SP_5D_I(0,2,is,Ns,site)] = v2_13;
533  v2[IDX2_SP_5D_R(1,2,is,Ns,site)] = v2_23;
534  v2[IDX2_SP_5D_I(1,2,is,Ns,site)] = v2_33;
535  v2[IDX2_SP_5D_R(2,2,is,Ns,site)] = v2_43;
536  v2[IDX2_SP_5D_I(2,2,is,Ns,site)] = v2_53;
537 
538  v2[IDX2_SP_5D_R(0,3,is,Ns,site)] = v2_04;
539  v2[IDX2_SP_5D_I(0,3,is,Ns,site)] = v2_14;
540  v2[IDX2_SP_5D_R(1,3,is,Ns,site)] = v2_24;
541  v2[IDX2_SP_5D_I(1,3,is,Ns,site)] = v2_34;
542  v2[IDX2_SP_5D_R(2,3,is,Ns,site)] = v2_44;
543  v2[IDX2_SP_5D_I(2,3,is,Ns,site)] = v2_54;
544  }
545 
546  } // site loop
547  } // acc parallel
548  } // acc data
549 
550 }
551 
552 //====================================================================
553 void mult_domainwall_5din_eo_hop1_dirac(
554  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
555  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
556  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
557  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
558  real_t *RESTRICT up, real_t *RESTRICT wp,
559  int Ns, int *bc, int *Nsize, int *do_comm, int ieo, int jeo,
560  int jgm5)
561 {
562  int Nx = Nsize[0];
563  int Ny = Nsize[1];
564  int Nz = Nsize[2];
565  int Nt = Nsize[3];
566  int Nst = Nx * Ny * Nz * Nt;
567  int Nst_pad = CEIL_NWP(Nst);
568 
569  int Nin5 = NVCD * Ns;
570  int Nin5bd = NVC * ND2 * Ns;
571 
572  int size = Nin5 * Nst_pad;
573  int size_u = NDF * Nst_pad * 2 * NDIM;
574 
575  int size_bx = Nin5bd * CEIL_NWP((Ny * Nz * Nt + 1)/2);
576  int size_by = Nin5bd * CEIL_NWP(Nst/Ny);
577  int size_bz = Nin5bd * CEIL_NWP(Nst/Nz);
578  int size_bt = Nin5bd * CEIL_NWP(Nst/Nt);
579 
580 #pragma acc data present(up[0:size_u], wp[0:size], \
581  buf_xp[0:size_bx], buf_xm[0:size_bx], \
582  buf_yp[0:size_by], buf_ym[0:size_by], \
583  buf_zp[0:size_bz], buf_zm[0:size_bz], \
584  buf_tp[0:size_bt], buf_tm[0:size_bt] ) \
585  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
586  bc[0:NDIM], do_comm[0:4], ieo, jeo, jgm5)
587  {
588 
589  if (do_comm[0] > 0) {
590  int idir = 0;
591  int Nyzt = Ny * Nz * Nt;
592  int Nyzt_pad = CEIL_NWP(Nyzt);
593 
594 #pragma acc parallel async \
595  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
596  {
597 
598 #pragma acc loop gang worker vector
599  for (int iyzt = 0; iyzt < Nyzt_pad; ++iyzt) {
600  if(iyzt < Nyzt){
601 
602  int iy = iyzt % Ny;
603  int iz = (iyzt/Ny) % Nz;
604  int it = iyzt/(Ny * Nz);
605  int keo = (jeo + iy + iz + it) % 2;
606 
607  if(keo == 1){
608  int ix = 0;
609  int iyzt2 = iyzt/2;
610  int site = ix + Nx * iyzt;
611  real_t bc2 = bc[0];
612  for (int is = 0; is < Ns; ++is) {
613  real_t wt[NVCD], vt[NVC * ND2];
614  if(jgm5 == 0){
615  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
616  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
617  }
618  }else{
619  for(int ivc = 0; ivc < NVC; ++ivc){
620  wt[ivc+ID1] = wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), site)];
621  wt[ivc+ID2] = wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), site)];
622  wt[ivc+ID3] = wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), site)];
623  wt[ivc+ID4] = wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), site)];
624  }
625  }
626  mult_wilson_xp1(vt, wt);
627  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
628  buf_xp[IDX2(Nin5bd, (ivcd + NVCD2 * is), iyzt2)] = bc2 * vt[ivcd];
629  }
630  }
631  }
632 
633  if(keo == 0){
634  int iyzt2 = iyzt/2;
635  int ix = Nx-1;
636  int site = ix + Nx * iyzt;
637  real_t ut[NDF];
638  load_u(ut, up, site + Nst_pad * (1-ieo + 2*idir));
639  for (int is = 0; is < Ns; ++is) {
640  real_t wt[NVCD], vt[NVC * ND2];
641  if(jgm5 == 0){
642  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
643  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
644  }
645  }else{
646  for(int ivc = 0; ivc < NVC; ++ivc){
647  wt[ivc+ID1] = wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), site)];
648  wt[ivc+ID2] = wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), site)];
649  wt[ivc+ID3] = wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), site)];
650  wt[ivc+ID4] = wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), site)];
651  }
652  }
653  mult_wilson_xm1(vt, ut, wt);
654  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
655  buf_xm[IDX2(Nin5bd, (ivcd + NVCD2 * is), iyzt2)] = vt[ivcd];
656  }
657  }
658  }
659  }
660  }
661  }
662  } // do_comm[0]
663 
664  if (do_comm[1] > 0) {
665  int idir = 1;
666  int Nxzt = Nx * Nz * Nt;
667 
668 #pragma acc parallel async \
669  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
670  {
671  int Nxzt_pad = CEIL_NWP(Nxzt);
672 
673 #pragma acc loop gang worker vector
674  for (int ixzt = 0; ixzt < Nxzt_pad; ++ixzt) {
675  if(ixzt < Nxzt){
676 
677  int iy = 0;
678  int ix = ixzt % Nx;
679  int izt = ixzt / Nx;
680  int site = ix + Nx * (iy + Ny * izt);
681  real_t bc2 = bc[1];
682  for (int is = 0; is < Ns; ++is) {
683  real_t wt[NVCD], vt[NVC * ND2];
684  if(jgm5 == 0){
685  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
686  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
687  }
688  }else{
689  for(int ivc = 0; ivc < NVC; ++ivc){
690  wt[ivc+ID1] = wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), site)];
691  wt[ivc+ID2] = wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), site)];
692  wt[ivc+ID3] = wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), site)];
693  wt[ivc+ID4] = wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), site)];
694  }
695  }
696  mult_wilson_yp1(vt, wt);
697  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
698  buf_yp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixzt)] = bc2 * vt[ivcd];
699  }
700  }
701 
702  iy = Ny-1;
703  site = ix + Nx * (iy + Ny * izt);
704  real_t ut[NDF];
705  load_u(ut, up, site + Nst_pad * (1-ieo + 2*idir));
706  for (int is = 0; is < Ns; ++is) {
707  real_t wt[NVCD], vt[NVC * ND2];
708  if(jgm5 == 0){
709  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
710  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
711  }
712  }else{
713  for(int ivc = 0; ivc < NVC; ++ivc){
714  wt[ivc+ID1] = wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), site)];
715  wt[ivc+ID2] = wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), site)];
716  wt[ivc+ID3] = wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), site)];
717  wt[ivc+ID4] = wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), site)];
718  }
719  }
720  mult_wilson_ym1(vt, ut, wt);
721  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
722  buf_ym[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixzt)] = vt[ivcd];
723  }
724  }
725  }
726  }
727  }
728  } // do_comm[1]
729 
730  if (do_comm[2] > 0) {
731  int idir = 2;
732  int Nxy = Nx * Ny;
733  int Nxyt = Nx * Ny * Nt;
734 
735 #pragma acc parallel async \
736  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
737  {
738  int Nxyt_pad = CEIL_NWP(Nxyt);
739 
740 #pragma acc loop gang worker vector
741  for (int ixyt = 0; ixyt < Nxyt_pad; ++ixyt) {
742  if(ixyt < Nxyt){
743 
744  int iz = 0;
745  int ixy = ixyt % Nxy;
746  int it = ixyt / Nxy;
747  int site = ixy + Nxy * (iz + Nz * it);
748  real_t bc2 = bc[2];
749  for (int is = 0; is < Ns; ++is) {
750  real_t wt[NVCD], vt[NVC * ND2];
751  if(jgm5 == 0){
752  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
753  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
754  }
755  }else{
756  for(int ivc = 0; ivc < NVC; ++ivc){
757  wt[ivc+ID1] = wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), site)];
758  wt[ivc+ID2] = wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), site)];
759  wt[ivc+ID3] = wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), site)];
760  wt[ivc+ID4] = wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), site)];
761  }
762  }
763  mult_wilson_zp1(vt, wt);
764  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
765  buf_zp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyt)] = bc2 * vt[ivcd];
766  }
767  }
768 
769  iz = Nz-1;
770  site = ixy + Nxy * (iz + Nz * it);
771  real_t ut[NDF];
772  load_u(ut, up, site + Nst_pad * (1-ieo + 2*idir));
773  for (int is = 0; is < Ns; ++is) {
774  real_t wt[NVCD], vt[NVC * ND2];
775  if(jgm5 == 0){
776  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
777  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
778  }
779  }else{
780  for(int ivc = 0; ivc < NVC; ++ivc){
781  wt[ivc+ID1] = wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), site)];
782  wt[ivc+ID2] = wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), site)];
783  wt[ivc+ID3] = wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), site)];
784  wt[ivc+ID4] = wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), site)];
785  }
786  }
787  mult_wilson_zm1(vt, ut, wt);
788  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
789  buf_zm[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyt)] = vt[ivcd];
790  }
791  }
792  }
793  }
794 
795  }
796  } // do_comm[2]
797 
798  if (do_comm[3] > 0) {
799  int idir = 3;
800  int Nxyz = Nx * Ny * Nz;
801 
802 #pragma acc parallel async \
803  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
804  {
805  int Nxyz_pad = CEIL_NWP(Nxyz);
806 
807 #pragma acc loop gang worker vector
808  for (int ixyz = 0; ixyz < Nxyz_pad; ++ixyz) {
809  if(ixyz < Nxyz){
810 
811  int it = 0;
812  int site = ixyz + Nxyz * it;
813  real_t bc2 = bc[3];
814  for (int is = 0; is < Ns; ++is) {
815  real_t wt[NVCD], vt[NVC * ND2];
816  if(jgm5 == 0){
817  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
818  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
819  }
820  }else{
821  for(int ivc = 0; ivc < NVC; ++ivc){
822  wt[ivc+ID1] = wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), site)];
823  wt[ivc+ID2] = wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), site)];
824  wt[ivc+ID3] = wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), site)];
825  wt[ivc+ID4] = wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), site)];
826  }
827  }
828  mult_wilson_tp1_dirac(vt, wt);
829  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
830  buf_tp[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyz)] = bc2 * vt[ivcd];
831  }
832  }
833 
834  it = Nt-1;
835  site = ixyz + Nxyz * it;
836  real_t ut[NDF];
837  load_u(ut, up, site + Nst_pad * (1-ieo + 2*idir));
838  for (int is = 0; is < Ns; ++is) {
839  real_t wt[NVCD], vt[NVC * ND2];
840  if(jgm5 == 0){
841  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
842  wt[ivcd] = wp[IDX2(Nin5, (ivcd + NVCD * is), site)];
843  }
844  }else{
845  for(int ivc = 0; ivc < NVC; ++ivc){
846  wt[ivc+ID1] = wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), site)];
847  wt[ivc+ID2] = wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), site)];
848  wt[ivc+ID3] = wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), site)];
849  wt[ivc+ID4] = wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), site)];
850  }
851  }
852  mult_wilson_tm1_dirac(vt, ut, wt);
853  for(int ivcd = 0; ivcd < NVC * ND2; ++ivcd){
854  buf_tm[IDX2(Nin5bd, (ivcd + NVCD2 * is), ixyz)] = vt[ivcd];
855  }
856  }
857  }
858  }
859 
860  }
861  } // do_comm[3]
862 
863 #pragma acc wait
864 
865  } // acc data
866 
867  if(do_comm[0] > 0){
868  #pragma acc update async host (buf_xp[0:size_bx])
869  #pragma acc update async host (buf_xm[0:size_bx])
870  }
871  if(do_comm[1] > 0){
872  #pragma acc update async host (buf_yp[0:size_by])
873  #pragma acc update async host (buf_ym[0:size_by])
874  }
875  if(do_comm[2] > 0){
876  #pragma acc update async host (buf_zp[0:size_bz])
877  #pragma acc update async host (buf_zm[0:size_bz])
878  }
879  if(do_comm[3] > 0){
880  #pragma acc update async host (buf_tp[0:size_bt])
881  #pragma acc update async host (buf_tm[0:size_bt])
882  }
883 
884  #pragma acc wait
885 
886 }
887 
888 //====================================================================
889 void mult_domainwall_5din_eo_hop2_dirac(
890  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
891  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
892  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
893  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
894  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
895  int Ns, int *bc, int *Nsize, int *do_comm, int ieo, int jeo)
896 {
897  int Nx = Nsize[0];
898  int Ny = Nsize[1];
899  int Nz = Nsize[2];
900  int Nt = Nsize[3];
901  int Nst = Nx * Ny * Nz * Nt;
902  int Nst_pad = CEIL_NWP(Nst);
903 
904  int Nin5 = NVCD * Ns;
905  int Nin5bd = (NVCD/2) * Ns;
906 
907  int size = Nin5 * Nst_pad;
908  int size_u = NDF * Nst_pad * 2 * NDIM;
909  int size_bx = Nin5bd * CEIL_NWP((Ny * Nz * Nt + 1)/2);
910  int size_by = Nin5bd * CEIL_NWP(Nx * Nz * Nt);
911  int size_bz = Nin5bd * CEIL_NWP(Nx * Ny * Nt);
912  int size_bt = Nin5bd * CEIL_NWP(Nx * Ny * Nz);
913 
914  if(do_comm[0] > 0){
915  #pragma acc update async device (buf_xp[0:size_bx])
916  #pragma acc update async device (buf_xm[0:size_bx])
917  }
918  if(do_comm[1] > 0){
919  #pragma acc update async device (buf_yp[0:size_by])
920  #pragma acc update async device (buf_ym[0:size_by])
921  }
922  if(do_comm[2] > 0){
923  #pragma acc update async device (buf_zp[0:size_bz])
924  #pragma acc update async device (buf_zm[0:size_bz])
925  }
926  if(do_comm[3] > 0){
927  #pragma acc update async device (buf_tp[0:size_bt])
928  #pragma acc update async device (buf_tm[0:size_bt])
929  }
930 
931 #pragma acc wait
932 
933 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
934  buf_yp[0:size_by], buf_ym[0:size_by], \
935  buf_zp[0:size_bz], buf_zm[0:size_bz], \
936  buf_tp[0:size_bt], buf_tm[0:size_bt], \
937  vp[0:size], up[0:size_u]) \
938  copyin(Nst, Nx, Ny, Nz, Nt, Nin5, Ns, \
939  bc[0:NDIM], do_comm[0:4], ieo, jeo)
940  {
941 
942 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
943  {
944  int Nxy = Nx * Ny;
945  int Nxyz = Nx * Ny * Nz;
946  int Nst_pad = CEIL_NWP(Nst);
947 
948 #pragma acc loop gang worker vector
949  for (int site = 0; site < Nst_pad; ++site) {
950  if(site < Nst){
951 
952  int ix = site % Nx;
953  int iyzt = site / Nx;
954  int ixy = site % Nxy;
955  int iy = iyzt % Ny;
956  int izt = site / Nxy;
957  int iz = izt % Nz;
958  int it = izt / Nz;
959  int ixyz = site % Nxyz;
960  int keo = (jeo + iy + iz + it) % 2;
961 
962  int idir;
963 
964  for(int is = 0; is < Ns; ++is){
965 
966  real_t vL[NVCD];
967 
968  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
969  vL[ivcd] = 0.0;
970  }
971 
972  int opr_any = 0;
973 
974  idir = 0;
975  if (do_comm[idir] > 0) {
976 
977  if(ix == Nx-1 && keo == 1){
978  real_t ut[NDF];
979  load_u(ut, up, site + Nst_pad * (ieo + 2*idir));
980  real_t wt[NVC * ND2];
981  int iyzt2 = iyzt/2;
982  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
983  wt[ivcd] = buf_xp[IDX2(Nin5bd, ivcd + NVCD2 * is, iyzt2)];
984  }
985  mult_wilson_xp2(vL, ut, wt);
986  ++opr_any;
987  }
988 
989  if(ix == 0 && keo == 0){
990  real_t bc2 = bc[0];
991  int iyzt2 = iyzt/2;
992  real_t wt[NVC * ND2];
993  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
994  wt[ivcd] = bc2 * buf_xm[IDX2(Nin5bd, ivcd + NVCD2 * is, iyzt2)];
995  }
996  mult_wilson_xm2(vL, wt);
997  ++opr_any;
998  }
999  }
1000 
1001  idir = 1;
1002  if (do_comm[idir] > 0) {
1003  int ixzt = ix + Nx * izt;
1004 
1005  if (iy == Ny-1) {
1006  real_t ut[NDF];
1007  load_u(ut, up, site + Nst_pad * (ieo + 2*idir));
1008  real_t wt[NVC * ND2];
1009  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1010  wt[ivcd] = buf_yp[IDX2(Nin5bd, ivcd + NVCD2 * is, ixzt)];
1011  }
1012  mult_wilson_yp2(vL, ut, wt);
1013  ++opr_any;
1014  }
1015 
1016  if (iy == 0) {
1017  real_t bc2 = bc[1];
1018  real_t wt[NVC * ND2];
1019  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1020  wt[ivcd] = bc2 * buf_ym[IDX2(Nin5bd, ivcd + NVCD2 * is, ixzt)];
1021  }
1022  mult_wilson_ym2(vL, wt);
1023  ++opr_any;
1024  }
1025  }
1026 
1027  idir = 2;
1028  if (do_comm[idir] > 0) {
1029  int ixyt = ixy + Nxy * it;
1030  if (iz == Nz-1) {
1031  real_t ut[NDF];
1032  load_u(ut, up, site + Nst_pad * (ieo + 2*idir));
1033  real_t wt[NVC * ND2];
1034  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1035  wt[ivcd] = buf_zp[IDX2(Nin5bd, ivcd + NVCD2 * is, ixyt)];
1036  }
1037  mult_wilson_zp2(vL, ut, wt);
1038  ++opr_any;
1039  }
1040 
1041  if (iz == 0) {
1042  real_t bc2 = bc[2];
1043  real_t wt[NVC * ND2];
1044  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1045  wt[ivcd] = bc2 * buf_zm[IDX2(Nin5bd, ivcd + NVCD2 * is, ixyt)];
1046  }
1047  mult_wilson_zm2(vL, wt);
1048  ++opr_any;
1049  }
1050  }
1051 
1052  idir = 3;
1053  if (do_comm[idir] > 0) {
1054  if (it == Nt-1) {
1055  real_t ut[NDF];
1056  load_u(ut, up, site + Nst_pad * (ieo + 2*idir));
1057  real_t wt[NVC * ND2];
1058  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1059  wt[ivcd] = buf_tp[IDX2(Nin5bd, ivcd + NVCD2 * is, ixyz)];
1060  }
1061  mult_wilson_tp2_dirac(vL, ut, wt);
1062  ++opr_any;
1063  }
1064 
1065  if (it == 0) {
1066  real_t bc2 = bc[3];
1067  real_t wt[NVC * ND2];
1068  for(int ivcd = 0; ivcd < NVCD2; ++ivcd){
1069  wt[ivcd] = bc2 * buf_tm[IDX2(Nin5bd, ivcd + NVCD2 * is, ixyz)];
1070  }
1071  mult_wilson_tm2_dirac(vL, wt);
1072  ++opr_any;
1073  }
1074  }
1075 
1076  if (opr_any > 0) {
1077  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
1078  vp[IDX2(Nin5, ivcd + NVCD * is, site)] += vL[ivcd];
1079  }
1080  }
1081 
1082  } // is loop
1083 
1084  }
1085  } // site loop
1086 
1087  } // acc parallel
1088  } // acc data
1089 
1090 }
1091 
1092 
1093 #endif
1094 //============================================================END=====
u14
u14
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:100
v2_24
v2_24
Definition: mult_Clover_csw_chiral_openacc-inc.h:105
v2_22
v2_22
Definition: mult_Clover_csw_chiral_openacc-inc.h:91
u13
u13
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:99
mult_domainwall_5din_ee_5dirdag_dirac
void mult_domainwall_5din_ee_5dirdag_dirac(real_t *RESTRICT vp, real_t *RESTRICT wp, real_t mq, real_t M0, int Ns, real_t *b, real_t *c, real_t alpha, int *Nsize)
Definition: mult_Domainwall_5din_eo_openacc-inc.h:230
BridgeACC::mult_wilson_yp2
void mult_wilson_yp2(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:467
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
vt2_1
vt2_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:40
v2_04
v2_04
Definition: mult_Clover_csw_chiral_openacc-inc.h:103
iy
int iy
Definition: mult_Wilson_xyz_openacc-inc.h:239
ixy
int ixy
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:513
ID1
#define ID1
Definition: fopr_Wilson_impl_SU2-inc.h:18
wt
real_t wt[NVCD]
Definition: mult_Clover_csw_chiral_openacc-inc.h:9
izt
int izt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:262
u_0
u_0
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:46
NVCD
#define NVCD
Definition: define_params_SU3.h:20
v2_11
v2_11
Definition: mult_Clover_csw_chiral_openacc-inc.h:83
BridgeACC::mult_wilson_xm2
void mult_wilson_xm2(double *RESTRICT v2, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:282
mult_domainwall_5din_eo_hop1_dirac
void mult_domainwall_5din_eo_hop1_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 up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int ieo, int jeo, int jgm5)
Definition: mult_Domainwall_5din_eo_openacc-inc.h:553
mult_domainwall_5din_eo_hopb_dirac_5d
void mult_domainwall_5din_eo_hopb_dirac_5d(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, int Ns, int *bc, int *Nsize, int *do_comm, int ieo, int jeo, int jgm5)
Definition: mult_Domainwall_5din_eo_openacc-inc.h:450
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
mult_Domainwall_eo_xyz_openacc-inc.h
BridgeACC::mult_wilson_yp1
void mult_wilson_yp1(double *RESTRICT buf, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:414
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
vt4
real_t vt4
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
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
IDX2_SP_5D_I
#define IDX2_SP_5D_I(ic, id, is, Ns, ist)
Definition: define_index.h:62
v2_23
v2_23
Definition: mult_Clover_csw_chiral_openacc-inc.h:98
BridgeACC::mult_wilson_zp1
void mult_wilson_zp1(double *RESTRICT buf, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:836
u_3
u_3
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:49
BridgeACC::mult_wilson_tp2_dirac
void mult_wilson_tp2_dirac(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1311
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
BridgeACC::mult_wilson_tm2_dirac
void mult_wilson_tm2_dirac(double *RESTRICT v2, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1522
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
iyzt
int iyzt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:14
ID4
#define ID4
Definition: fopr_Wilson_impl_SU2-inc.h:21
mult_domainwall_5din_ee_5dir_dirac
void mult_domainwall_5din_ee_5dir_dirac(real_t *RESTRICT vp, real_t *RESTRICT wp, real_t mq, real_t M0, int Ns, real_t *b, real_t *c, real_t alpha, int *Nsize)
Definition: mult_Domainwall_5din_eo_openacc-inc.h:14
ix
int ix
Definition: mult_Wilson_xyz_openacc-inc.h:16
vt1_0
vt1_0
Definition: mult_Wilson_eo_t_chiral_openacc-inc.h:18
idir
idir
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:264
u10
u10
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:73
vt1
real_t vt1
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
BridgeACC::mult_wilson_xm1
void mult_wilson_xm1(double *RESTRICT buf, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:209
BridgeACC::mult_wilson_ym2
void mult_wilson_ym2(double *RESTRICT v2, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:696
bc2
int bc2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:62
wt1r
wt1r
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:53
wt2i
wt2i
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:59
AIndex_eo_accel::idx
int idx(const int in, const int Nin, const int ist, const int leo, const int Nvol2, const int ex)
Definition: aindex_eo.h:28
mult_domainwall_5din_eo_hop2_dirac
void mult_domainwall_5din_eo_hop2_dirac(real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp, 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, int Ns, int *bc, int *Nsize, int *do_comm, int ieo, int jeo)
Definition: mult_Domainwall_5din_eo_openacc-inc.h:889
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
u17
u17
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:103
BridgeACC::mult_wilson_tm1_dirac
void mult_wilson_tm1_dirac(double *RESTRICT buf, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1450
it
int it
Definition: mult_Wilson_xyz_openacc-inc.h:461
v2_42
v2_42
Definition: mult_Clover_csw_chiral_openacc-inc.h:93
BridgeACC::mult_wilson_tp1_dirac
void mult_wilson_tp1_dirac(double *RESTRICT buf, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1261
BridgeACC::mult_wilson_xp2
void mult_wilson_xp2(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:64
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
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
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
IDX2
#define IDX2(nin, in, ist)
Definition: define_index.h:28
u12
u12
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:98
mult_domainwall_5din_eo_5dir_dirac
void mult_domainwall_5din_eo_5dir_dirac(real_t *RESTRICT yp, real_t *RESTRICT wp, real_t mq, real_t M0, int Ns, real_t *b, real_t *c, real_t alpha, int *Nsize)
Definition: mult_Domainwall_5din_eo_openacc-inc.h:122
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
BridgeACC::mult_wilson_zm2
void mult_wilson_zm2(double *RESTRICT v2, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1121
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
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
mult_domainwall_5din_eo_5dirdag_dirac
void mult_domainwall_5din_eo_5dirdag_dirac(real_t *RESTRICT vp, real_t *RESTRICT yp, real_t mq, real_t M0, int Ns, real_t *b, real_t *c, real_t alpha, int *Nsize)
Definition: mult_Domainwall_5din_eo_openacc-inc.h:346
load_u
void load_u(real_t *ut, real_t *up, int site)
Definition: mult_Wilson_inline_openacc-inc.h:26
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
v2_44
v2_44
Definition: mult_Clover_csw_chiral_openacc-inc.h:107
NWP
#define NWP
Definition: define_params.h:32
BridgeACC::mult_wilson_zp2
void mult_wilson_zp2(double *RESTRICT v2, double *RESTRICT u, double *RESTRICT buf, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:891
iz
int iz
Definition: mult_Wilson_xyz_openacc-inc.h:462
BridgeACC::mult_wilson_ym1
void mult_wilson_ym1(double *RESTRICT buf, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:621
NDIM
#define NDIM
Definition: contract_4spinor.cpp:18
ixyz
int ixyz
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:13
u15
u15
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:101
BridgeACC::mult_wilson_zm1
void mult_wilson_zm1(double *RESTRICT buf, double *RESTRICT u, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:1045
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
BridgeACC::mult_wilson_xp1
void mult_wilson_xp1(double *RESTRICT buf, double *RESTRICT v1, int *Nsize, int *bc, int Nc)
Definition: mult_Wilson_dir_openacc-inc.h:14
v2_43
v2_43
Definition: mult_Clover_csw_chiral_openacc-inc.h:100
vt2_0
vt2_0
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:39
vt2
real_t vt2
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
ND2
#define ND2
Definition: define_params_SU3.h:18
IDX2_SP_5D_R
#define IDX2_SP_5D_R(ic, id, is, Ns, ist)
Definition: define_index.h:61
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
vt3
real_t vt3
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
vt1_4
vt1_4
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:36
vt1_1
vt1_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:33
mult_Domainwall_eo_t_dirac_openacc-inc.h