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