Bridge++  Ver.2.1.3
mult_Staggered_openacc-inc.h
Go to the documentation of this file.
1 
9 // This code explicitly assumes SU(3) gauge group.
10 
11 #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)
12 #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)
13 
14 //====================================================================
15 void mult_staggered_D(real_t *RESTRICT v2, real_t *RESTRICT u_up,
16  real_t *RESTRICT u_dn, real_t *RESTRICT v1,
17  real_t mq, int *bc, int *Nsize, int jdag)
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 size = NVC * Nst_pad;
27  int size_u = NDF * Nst_pad * 4;
28 
29 #pragma acc data present(v2[0:size], v1[0:size], \
30  u_up[0:size_u], u_dn[0:size_u]), \
31  copyin(bc[0:4], Nx, Ny, Nz, Nt, Nst, Nst_pad, mq, jdag)
32 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
33  {
34 #pragma acc loop gang worker vector
35  for(int ist = 0; ist < Nst; ++ist){
36  int ix = ist % Nx;
37  int iy = (ist/Nx) % Ny;
38  int iz = (ist/(Nx * Ny)) % Nz;
39  int it = ist/(Nx * Ny * Nz);
40 
41  int idir, nei, istg, neig;
42  real_t bc2;
43  real_t vt0, vt1, vt2, vt3, vt4, vt5;
44  real_t ut0, ut1, ut2, ut3, ut4, ut5;
45  real_t xt0, xt1, xt2, xt3, xt4, xt5;
46  real_t wtr, wti;
47 
48  xt0 = 0.0;
49  xt1 = 0.0;
50  xt2 = 0.0;
51  xt3 = 0.0;
52  xt4 = 0.0;
53  xt5 = 0.0;
54 
55 
56  idir = 0;
57 
58  int iyzt = ist/Nx;
59  nei = ((ix+1) % Nx) + Nx * iyzt;
60  istg = ist + Nst_pad * idir;
61 
62  bc2 = 1.0;
63  if(ix == Nx-1) bc2 = bc[0];
64 
65 #include "inc/mult_Staggered_uvup_openacc-inc.h"
66 
67  nei = ix-1 + Nx * iyzt;
68  if(ix == 0) nei = Nx-1 + Nx * iyzt;
69  neig = nei + Nst_pad * idir;
70  bc2 = 1.0;
71  if(ix == 0) bc2 = bc[0];
72 
73 #include "inc/mult_Staggered_uvdn_openacc-inc.h"
74 
75 
76  idir = 1;
77  int izt = ist/(Nx*Ny);
78  int iy2 = (iy+1) % Ny;
79  nei = ix + Nx * (iy2 + Ny * izt);
80  istg = ist + Nst_pad * idir;
81  bc2 = 1.0;
82  if(iy == Ny-1) bc2 = bc[idir];
83 
84 #include "inc/mult_Staggered_uvup_openacc-inc.h"
85 
86  iy2 = (iy-1+Ny) % Ny;
87  nei = ix + Nx * (iy2 + Ny * izt);
88  neig = nei + Nst_pad * idir;
89 
90  bc2 = 1.0;
91  if(iy == 0) bc2 = bc[1];
92 
93 #include "inc/mult_Staggered_uvdn_openacc-inc.h"
94 
95 
96  idir = 2;
97 
98  int Nxy = Nx * Ny;
99  int ixy = ix + Nx * iy;
100  int iz2 = (iz+1) % Nz;
101  nei = ixy + Nxy * (iz2 + Nz * it);
102  istg = ist + Nst_pad * idir;
103  bc2 = 1.0;
104  if(iz == Nz-1) bc2 = bc[idir];
105 
106 #include "inc/mult_Staggered_uvup_openacc-inc.h"
107 
108  iz2 = (iz-1+Nz) % Nz;
109  nei = ixy + Nxy * (iz2 + Nz * it);
110  neig = nei + Nst_pad * idir;
111  bc2 = 1.0;
112  if(iz == 0) bc2 = bc[idir];
113 
114 #include "inc/mult_Staggered_uvdn_openacc-inc.h"
115 
116 
117  idir = 3;
118  int Nxyz = Nx * Ny * Nz;
119  int ixyz = ix + Nx * (iy + Ny * iz);
120  int it2 = (it+1) % Nt;
121  nei = ixyz + Nxyz * it2;
122  istg = ist + Nst_pad * idir;
123  bc2 = 1.0;
124  if(it == Nt-1) bc2 = bc[idir];
125 
126 #include "inc/mult_Staggered_uvup_openacc-inc.h"
127 
128  it2 = (it-1+Nt) % Nt;
129  nei = ixyz + Nxyz * it2;
130  neig = nei + Nst_pad * idir;
131  bc2 = 1.0;
132  if(it == 0) bc2 = bc[idir];
133 
134 #include "inc/mult_Staggered_uvdn_openacc-inc.h"
135 
136  real_t fac = real_t(jdag) * 0.5;
137  v2[IDX2_1SP_R(0, ist)] = mq * v1[IDX2_1SP_R(0, ist)] + fac * xt0;
138  v2[IDX2_1SP_I(0, ist)] = mq * v1[IDX2_1SP_I(0, ist)] + fac * xt1;
139  v2[IDX2_1SP_R(1, ist)] = mq * v1[IDX2_1SP_R(1, ist)] + fac * xt2;
140  v2[IDX2_1SP_I(1, ist)] = mq * v1[IDX2_1SP_I(1, ist)] + fac * xt3;
141  v2[IDX2_1SP_R(2, ist)] = mq * v1[IDX2_1SP_R(2, ist)] + fac * xt4;
142  v2[IDX2_1SP_I(2, ist)] = mq * v1[IDX2_1SP_I(2, ist)] + fac * xt5;
143 
144  }
145 
146  }
147 
148 }
149 
150 //====================================================================
151 void mult_staggered_1(real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
152  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
153  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
154  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
155  real_t *RESTRICT u_dn, real_t *RESTRICT v1,
156  int *bc, int *Nsize, int *do_comm)
157 {
158  int Nx = Nsize[0];
159  int Ny = Nsize[1];
160  int Nz = Nsize[2];
161  int Nt = Nsize[3];
162  int Nst = Nx * Ny * Nz * Nt;
163  int Nst_pad = CEIL_NWP(Nst);
164 
165  int size = NVC * Nst_pad;
166  int size_u = NDF * Nst_pad * NDIM;
167  int size_bx = NVC * CEIL_NWP(Ny * Nz * Nt);
168  int size_by = NVC * CEIL_NWP(Nx * Nz * Nt);
169  int size_bz = NVC * CEIL_NWP(Nx * Ny * Nt);
170  int size_bt = NVC * CEIL_NWP(Nx * Ny * Nz);
171 
172 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
173  buf_yp[0:size_by], buf_ym[0:size_by], \
174  buf_zp[0:size_bz], buf_zm[0:size_bz], \
175  buf_tp[0:size_bt], buf_tm[0:size_bt], \
176  u_dn[0:size_u], v1[0:size]) \
177  copyin(bc[0:4], do_comm[0:4], Nx, Ny, Nz, Nt, Nst, Nst_pad)
178  {
179 
180  if(do_comm[0] > 0){
181 
182 #pragma acc parallel async \
183  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
184  {
185  int Nyzt = Ny * Nz * Nt;
186 
187 #pragma acc loop gang worker vector
188  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
189  int ix = 0;
190  int ist = ix + Nx * iyzt;
191  real_t bc2 = bc[0];
192 
193  real_t *wt = buf_xp;
194 
195  for(int ic = 0; ic < NC; ++ic){
196  wt[IDX2_1SP_R(ic, iyzt)] = bc2 * v1[IDX2_1SP_R(ic, ist)];
197  wt[IDX2_1SP_I(ic, iyzt)] = bc2 * v1[IDX2_1SP_I(ic, ist)];
198  }
199  }
200 
201  }
202 
203 #pragma acc parallel async \
204  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
205  {
206  int Nyzt = Ny * Nz * Nt;
207 
208 #pragma acc loop gang worker vector
209  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
210  int ix = Nx-1;
211  int ist = ix + Nx * iyzt;
212  int istu = ist;
213  real_t *wt = buf_xm;
214  int ibf = iyzt;
215 
216 #include "inc/mult_Staggered_uvdn1_openacc-inc.h"
217 
218  }
219  }
220 
221  } // do_comm[0]
222 
223 
224  // idir = 1;
225  if(do_comm[1] > 0){
226 
227 #pragma acc parallel async \
228  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
229  {
230  int Nzt = Nz * Nt;
231 
232 #pragma acc loop gang worker vector collapse(2)
233  for(int izt = 0; izt < Nzt; ++izt){
234  for(int ix = 0; ix < Nx; ++ix){
235  int iy = 0;
236  int ist = ix + Nx * (iy + Ny * izt);
237  int ixzt = ix + Nx * izt;
238  real_t bc2 = bc[1];
239  real_t *wt = buf_yp;
240  for(int ic = 0; ic < NC; ++ic){
241  wt[IDX2_1SP_R(ic, ixzt)] = bc2 * v1[IDX2_1SP_R(ic, ist)];
242  wt[IDX2_1SP_I(ic, ixzt)] = bc2 * v1[IDX2_1SP_I(ic, ist)];
243  }
244  }
245  }
246 
247  }
248 
249 #pragma acc parallel async \
250  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
251  {
252  int Nzt = Nz * Nt;
253 
254 #pragma acc loop gang worker vector collapse(2)
255  for(int izt = 0; izt < Nzt; ++izt){
256  for(int ix = 0; ix < Nx; ++ix){
257  int iy = Ny-1;
258  int ist = ix + Nx * (iy + Ny*izt);
259  int istu = ist + Nst_pad;
260 
261  int ibf = ix + Nx * izt;
262  real_t *wt = buf_ym;
263 
264 #include "inc/mult_Staggered_uvdn1_openacc-inc.h"
265 
266  }
267  }
268 
269  }
270 
271  } // do_comm[1]
272 
273  // idir = 2;
274  if(do_comm[2] > 0){
275 
276 #pragma acc parallel async \
277  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
278  {
279  int Nxy = Nx * Ny;
280 
281 #pragma acc loop gang worker vector collapse(2)
282  for(int it = 0; it < Nt; ++it){
283  for(int ixy = 0; ixy < Nxy; ++ixy){
284  int iz = 0;
285  int ist = ixy + Nxy * (iz + Nz * it);
286  int ixyt = ixy + Nxy * it;
287 
288  real_t bc2 = bc[2];
289 
290  real_t *wt = buf_zp;
291 
292  for(int ic = 0; ic < NC; ++ic){
293  wt[IDX2_1SP_R(ic, ixyt)] = bc2 * v1[IDX2_1SP_R(ic, ist)];
294  wt[IDX2_1SP_I(ic, ixyt)] = bc2 * v1[IDX2_1SP_I(ic, ist)];
295  }
296 
297  }
298  }
299 
300  }
301 
302 #pragma acc parallel async \
303  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
304  {
305  int Nxy = Nx * Ny;
306 
307 #pragma acc loop gang worker vector collapse(2)
308  for(int it = 0; it < Nt; ++it){
309  for(int ixy = 0; ixy < Nxy; ++ixy){
310  int iz = Nz-1;
311  int ist = ixy + Nxy * (iz + Nz * it);
312  int istu = ist + Nst_pad * 2;
313 
314  real_t *wt = buf_zm;
315  int ibf = ixy + Nxy * it;
316 
317 #include "inc/mult_Staggered_uvdn1_openacc-inc.h"
318 
319  }
320  }
321 
322  }
323 
324  } // do_comm[2]
325 
326  // idir = 3;
327  if(do_comm[3] > 0){
328 
329 #pragma acc parallel async \
330  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
331  {
332  int Nxyz = Nx * Ny * Nz;
333 
334 #pragma acc loop gang worker vector
335  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
336  int it = 0;
337  int ist = ixyz + Nxyz * it;
338  real_t bc2 = bc[3];
339 
340  real_t *wt = buf_tp;
341  for(int ic = 0; ic < NC; ++ic){
342  wt[IDX2_1SP_R(ic, ixyz)] = bc2 * v1[IDX2_1SP_R(ic, ist)];
343  wt[IDX2_1SP_I(ic, ixyz)] = bc2 * v1[IDX2_1SP_I(ic, ist)];
344  }
345 
346  }
347 
348  }
349 
350 #pragma acc parallel async \
351  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
352  {
353  int Nxyz = Nx * Ny * Nz;
354 
355 #pragma acc loop gang worker vector
356  for(int ixyz = 0; ixyz < Nxyz; ++ixyz){
357  int it = Nt-1;
358  int ist = ixyz + Nxyz * it;
359  int istu = ist + Nst_pad * 3;
360 
361  int ibf = ixyz;
362  real_t *wt = buf_tm;
363 
364 #include "inc/mult_Staggered_uvdn1_openacc-inc.h"
365 
366  }
367 
368  }
369 
370  } // do_comm[3]
371 
372  #pragma acc wait
373 
374  } // acc data
375 
376  if(do_comm[0] > 0){
377 #pragma acc update async host (buf_xp[0:size_bx])
378 #pragma acc update async host (buf_xm[0:size_bx])
379  }
380  if(do_comm[1] > 0){
381 #pragma acc update async host (buf_yp[0:size_by])
382 #pragma acc update async host (buf_ym[0:size_by])
383  }
384  if(do_comm[2] > 0){
385 #pragma acc update async host (buf_zp[0:size_bz])
386 #pragma acc update async host (buf_zm[0:size_bz])
387  }
388  if(do_comm[3] > 0){
389 #pragma acc update async host (buf_tp[0:size_bt])
390 #pragma acc update async host (buf_tm[0:size_bt])
391  }
392 
393  #pragma acc wait
394 
395 }
396 
397 //====================================================================
398 void mult_staggered_2(real_t *RESTRICT v2, real_t *RESTRICT u_up,
399  real_t *RESTRICT buf_xp, real_t *RESTRICT buf_xm,
400  real_t *RESTRICT buf_yp, real_t *RESTRICT buf_ym,
401  real_t *RESTRICT buf_zp, real_t *RESTRICT buf_zm,
402  real_t *RESTRICT buf_tp, real_t *RESTRICT buf_tm,
403  real_t mq, int *bc,
404  int *Nsize, int *do_comm, int jdag)
405 {
406  int Nx = Nsize[0];
407  int Ny = Nsize[1];
408  int Nz = Nsize[2];
409  int Nt = Nsize[3];
410  int Nst = Nx * Ny * Nz * Nt;
411  int Nst_pad = CEIL_NWP(Nst);
412 
413  int size = NVC * Nst_pad;
414  int size_u = NDF * Nst_pad * NDIM;
415  int size_bx = NVC * CEIL_NWP(Ny * Nz * Nt);
416  int size_by = NVC * CEIL_NWP(Nx * Nz * Nt);
417  int size_bz = NVC * CEIL_NWP(Nx * Ny * Nt);
418  int size_bt = NVC * CEIL_NWP(Nx * Ny * Nz);
419 
420  if(do_comm[0] > 0){
421 #pragma acc update async device (buf_xp[0:size_bx])
422 #pragma acc update async device (buf_xm[0:size_bx])
423  }
424  if(do_comm[1] > 0){
425 #pragma acc update async device (buf_yp[0:size_by])
426 #pragma acc update async device (buf_ym[0:size_by])
427  }
428  if(do_comm[2] > 0){
429 #pragma acc update async device (buf_zp[0:size_bz])
430 #pragma acc update async device (buf_zm[0:size_bz])
431  }
432  if(do_comm[3] > 0){
433 #pragma acc update async device (buf_tp[0:size_bt])
434 #pragma acc update async device (buf_tm[0:size_bt])
435  }
436 
437  #pragma acc wait
438 
439 
440 #pragma acc data present(buf_xp[0:size_bx], buf_xm[0:size_bx], \
441  buf_yp[0:size_by], buf_ym[0:size_by], \
442  buf_zp[0:size_bz], buf_zm[0:size_bz], \
443  buf_tp[0:size_bt], buf_tm[0:size_bt], \
444  v2[0:size], u_up[0:size_u]) \
445  copyin(bc[0:4], do_comm[0:4], mq, jdag, \
446  Nx, Ny, Nz, Nt, Nst, Nst_pad)
447  {
448 
449 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
450  {
451 
452 #pragma acc loop gang worker vector
453  for(int ist = 0; ist < Nst; ++ist){
454  int ix = ist % Nx;
455  int iy = (ist/Nx) % Ny;
456  int iz = (ist/(Nx*Ny)) % Nz;
457  int it = ist/(Nx*Ny*Nz);
458 
459  real_t xt0, xt1, xt2, xt3, xt4, xt5;
460  xt0 = 0.0;
461  xt1 = 0.0;
462  xt2 = 0.0;
463  xt3 = 0.0;
464  xt4 = 0.0;
465  xt5 = 0.0;
466 
467  if(do_comm[0] > 0){
468 
469  if(ix == Nx-1){
470  int iyzt = iy + Ny * (iz + Nz * it);
471  int istu = ist + Nst_pad * 0;
472  int ibuf = iyzt;
473 
474  real_t *buf = buf_xp;
475 
476 #include "inc/mult_Staggered_uvup2_openacc-inc.h"
477 
478  }
479 
480  if(ix == 0){
481 
482  int iyzt = iy + Ny * (iz + Nz * it);
483  int ibf = iyzt;
484  real_t bc2 = bc[0];
485  xt0 -= bc2 * buf_xm[IDX2_1SP_R(0, ibf)];
486  xt1 -= bc2 * buf_xm[IDX2_1SP_I(0, ibf)];
487  xt2 -= bc2 * buf_xm[IDX2_1SP_R(1, ibf)];
488  xt3 -= bc2 * buf_xm[IDX2_1SP_I(1, ibf)];
489  xt4 -= bc2 * buf_xm[IDX2_1SP_R(2, ibf)];
490  xt5 -= bc2 * buf_xm[IDX2_1SP_I(2, ibf)];
491  }
492 
493  }
494 
495 
496  if(do_comm[1] > 0){
497 
498  if(iy == Ny-1){
499  int ixzt = ix + Nx * (iz + Nz * it);
500  int istu = ist + Nst_pad * 1;
501 
502  int ibuf = ixzt;
503  real_t *buf = buf_yp;
504 
505 #include "inc/mult_Staggered_uvup2_openacc-inc.h"
506 
507  }
508 
509  if(iy == 0){
510  int ibf = ix + Nx * (iz + Nz * it);
511  real_t bc2 = bc[1];
512  xt0 -= bc2 * buf_ym[IDX2_1SP_R(0, ibf)];
513  xt1 -= bc2 * buf_ym[IDX2_1SP_I(0, ibf)];
514  xt2 -= bc2 * buf_ym[IDX2_1SP_R(1, ibf)];
515  xt3 -= bc2 * buf_ym[IDX2_1SP_I(1, ibf)];
516  xt4 -= bc2 * buf_ym[IDX2_1SP_R(2, ibf)];
517  xt5 -= bc2 * buf_ym[IDX2_1SP_I(2, ibf)];
518  }
519 
520  }
521 
522 
523  if(do_comm[2] > 0){
524 
525  if(iz == Nz-1){
526  int ixyt = ix + Nx * (iy + Ny * it);
527  int ibuf = ixyt;
528  int istu = ist + Nst_pad * 2;
529  real_t *buf = buf_zp;
530 
531 #include "inc/mult_Staggered_uvup2_openacc-inc.h"
532 
533  }
534 
535  if(iz == 0){
536  int ibf = ix + Nx * (iy + Ny * it);
537  real_t bc2 = bc[2];
538  xt0 -= bc2 * buf_zm[IDX2_1SP_R(0, ibf)];
539  xt1 -= bc2 * buf_zm[IDX2_1SP_I(0, ibf)];
540  xt2 -= bc2 * buf_zm[IDX2_1SP_R(1, ibf)];
541  xt3 -= bc2 * buf_zm[IDX2_1SP_I(1, ibf)];
542  xt4 -= bc2 * buf_zm[IDX2_1SP_R(2, ibf)];
543  xt5 -= bc2 * buf_zm[IDX2_1SP_I(2, ibf)];
544  }
545 
546  }
547 
548  if(do_comm[3] > 0){
549 
550  if(it == Nt-1){
551  int ibuf = ix + Nx * (iy + Ny * iz);
552  int istu = ist + Nst_pad * 3;
553 
554  real_t *buf = buf_tp;
555 
556 #include "inc/mult_Staggered_uvup2_openacc-inc.h"
557 
558  }
559 
560  if(it == 0){
561  int ibf = ix + Nx * (iy + Ny * iz);
562  real_t bc2 = bc[3];
563  xt0 -= bc2 * buf_tm[IDX2_1SP_R(0, ibf)];
564  xt1 -= bc2 * buf_tm[IDX2_1SP_I(0, ibf)];
565  xt2 -= bc2 * buf_tm[IDX2_1SP_R(1, ibf)];
566  xt3 -= bc2 * buf_tm[IDX2_1SP_I(1, ibf)];
567  xt4 -= bc2 * buf_tm[IDX2_1SP_R(2, ibf)];
568  xt5 -= bc2 * buf_tm[IDX2_1SP_I(2, ibf)];
569  }
570 
571  }
572 
573  // real_t fac = real_t(jdag) * 0.5/mq;
574  real_t fac = real_t(jdag) * 0.5;
575  v2[IDX2_1SP_R(0, ist)] += fac * xt0;
576  v2[IDX2_1SP_I(0, ist)] += fac * xt1;
577  v2[IDX2_1SP_R(1, ist)] += fac * xt2;
578  v2[IDX2_1SP_I(1, ist)] += fac * xt3;
579  v2[IDX2_1SP_R(2, ist)] += fac * xt4;
580  v2[IDX2_1SP_I(2, ist)] += fac * xt5;
581 
582  }
583 
584  }
585  }
586 
587 }
588 
589 
590 //============================================================END=====
mult_Staggered_uvup2_openacc-inc.h
xt1
xt1
Definition: mult_Staggered_uvdn_openacc-inc.h:30
xt5
xt5
Definition: mult_Staggered_uvdn_openacc-inc.h:62
IDX2_1SP_I
#define IDX2_1SP_I(ic, ist)
Definition: define_index.h:47
xt3
xt3
Definition: mult_Staggered_uvdn_openacc-inc.h:46
vt0
real_t vt0
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
iy
int iy
Definition: mult_Wilson_xyz_openacc-inc.h:239
ixy
int ixy
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:513
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
xt2
xt2
Definition: mult_Staggered_uvdn_openacc-inc.h:45
ut1
real_t ut1
Definition: mult_Staggered_uvdn1_openacc-inc.h:9
mult_Staggered_uvup_openacc-inc.h
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
vt4
real_t vt4
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
ut0
real_t ut0
Definition: mult_Staggered_uvdn1_openacc-inc.h:9
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
xt0
xt0
Definition: mult_Staggered_uvdn_openacc-inc.h:29
iyzt
int iyzt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:14
ix
int ix
Definition: mult_Wilson_xyz_openacc-inc.h:16
ut4
real_t ut4
Definition: mult_Staggered_uvdn1_openacc-inc.h:9
idir
idir
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:264
xt4
xt4
Definition: mult_Staggered_uvdn_openacc-inc.h:61
NC
#define NC
Definition: field_F_imp_SU2-inc.h:15
vt1
real_t vt1
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
bc2
int bc2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:62
mult_Staggered_uvdn_openacc-inc.h
mult_staggered_2
void mult_staggered_2(real_t *RESTRICT v2, real_t *RESTRICT u_up, 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 mq, int *bc, int *Nsize, int *do_comm, int jdag)
Definition: mult_Staggered_openacc-inc.h:398
ut5
real_t ut5
Definition: mult_Staggered_uvdn1_openacc-inc.h:9
wtr
real_t wtr
Definition: mult_Staggered_uvdn1_openacc-inc.h:10
vt5
real_t vt5
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
wti
real_t wti
Definition: mult_Staggered_uvdn1_openacc-inc.h:10
it
int it
Definition: mult_Wilson_xyz_openacc-inc.h:461
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
mult_staggered_1
void mult_staggered_1(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_dn, real_t *RESTRICT v1, int *bc, int *Nsize, int *do_comm)
Definition: mult_Staggered_openacc-inc.h:151
IDX2_1SP_R
#define IDX2_1SP_R(ic, ist)
Definition: define_index.h:46
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
mult_staggered_D
void mult_staggered_D(real_t *RESTRICT v2, real_t *RESTRICT u_up, real_t *RESTRICT u_dn, real_t *RESTRICT v1, real_t mq, int *bc, int *Nsize, int jdag)
Definition: mult_Staggered_openacc-inc.h:15
iz
int iz
Definition: mult_Wilson_xyz_openacc-inc.h:462
mult_Staggered_uvdn1_openacc-inc.h
NDIM
#define NDIM
Definition: contract_4spinor.cpp:18
ixyz
int ixyz
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:13
ut2
real_t ut2
Definition: mult_Staggered_uvdn1_openacc-inc.h:9
vt2
real_t vt2
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
vt3
real_t vt3
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
ut3
real_t ut3
Definition: mult_Staggered_uvdn1_openacc-inc.h:9