Bridge++  Ver.2.1.3
mult_Domainwall_5din_eo_4d_openacc-inc.h
Go to the documentation of this file.
1 
10 #ifndef MULT_DOMAINWALL_5DIN_EO_4D_ACC_INCLUDED
11 #define MULT_DOMAINWALL_5DIN_EO_4D_ACC_INCLUDED
12 
13 //====================================================================
14 void mult_domainwall_5din_eo_hopb_dirac_4d(
15  real_t *RESTRICT vp, real_t *RESTRICT up, real_t *RESTRICT wp,
16  int Ns, int *bc, int *Nsize, int *do_comm, int ieo, int jeo,
17  int jgm5)
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 
28  int size = Nin5 * Nst_pad;
29  int size_u = NDF * Nst_pad * 2 * NDIM;
30 
31 #pragma acc data present(vp[0:size], up[0:size_u], wp[0:size]) \
32  copyin(Nst, Nst_pad, Nx, Ny, Nz, Nt, Nin5, Ns, \
33  bc[0:NDIM], do_comm[0:NDIM], ieo, jeo, jgm5)
34  {
35 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
36  {
37 
38 #pragma acc loop gang worker vector
39  for (int site = 0; site < Nst_pad; ++site) {
40  if(site < Nst){
41  int Nxy = Nx * Ny;
42  int Nxyz = Nxy * Nz;
43 
44  int ix = site % Nx;
45  int iyzt = site / Nx;
46  int ixy = site % Nxy;
47  int iy = iyzt % Ny;
48  int izt = site / Nxy;
49  int iz = izt % Nz;
50  int it = izt / Nz;
51  int ixyz = site % Nxyz;
52 
53  int keo = (jeo + iy + iz + it) % 2;
54 
55  int idir;
56 
57  for(int is = 0; is < Ns; ++is){
58 
59  real_t vL[NVCD];
60 
61  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
62  vL[ivcd] = 0.0;
63  }
64 
65  idir = 0;
66 
67  // mult_xp
68  {
69  int ix2 = (ix + keo) % Nx;
70  int nei = ix2 + Nx * iyzt;
71  real_t bc2 = 1.0;
72  if(ix == Nx-1 && keo == 1) bc2 = bc[0];
73 
74  real_t ut[NDF];
75  load_u(ut, up, site + Nst_pad * (ieo + 2 * idir));
76 
77  real_t wt[NVCD];
78  if(jgm5 == 0){
79  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
80  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
81  }
82  }else{
83  for(int ivc = 0; ivc < NVC; ++ivc){
84  wt[ivc+ID1] = bc2 * wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), nei)];
85  wt[ivc+ID2] = bc2 * wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), nei)];
86  wt[ivc+ID3] = bc2 * wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), nei)];
87  wt[ivc+ID4] = bc2 * wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), nei)];
88  }
89  }
90  mult_wilson_xpb(vL, ut, wt);
91  }
92  // mult_xm
93  {
94  int ix2 = (ix - 1 + keo + Nx) % Nx;
95  int nei = ix2 + Nx * iyzt;
96  real_t bc2 = 1.0;
97  if(ix == 0 && keo == 0) bc2 = bc[0];
98 
99  real_t ut[NDF];
100  load_u(ut, up, nei + Nst_pad * (1-ieo + 2 * idir));
101 
102  real_t wt[NVCD];
103  if(jgm5 == 0){
104  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
105  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
106  }
107  }else{
108  for(int ivc = 0; ivc < NVC; ++ivc){
109  wt[ivc+ID1] = bc2 * wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), nei)];
110  wt[ivc+ID2] = bc2 * wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), nei)];
111  wt[ivc+ID3] = bc2 * wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), nei)];
112  wt[ivc+ID4] = bc2 * wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), nei)];
113  }
114  }
115  mult_wilson_xmb(vL, ut, wt);
116  }
117 
118  idir = 1;
119 
120  if ((iy < Ny-1) || (do_comm[idir] == 0)) {
121  int iy2 = (iy + 1) % Ny;
122  int nei = ix + Nx * (iy2 + Ny * izt);
123  real_t bc2 = 1.0;
124  if(iy == Ny-1) bc2 = bc[1];
125 
126  real_t ut[NDF];
127  load_u(ut, up, site + Nst_pad * (ieo + 2 * idir));
128 
129  real_t wt[NVCD];
130  if(jgm5 == 0){
131  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
132  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
133  }
134  }else{
135  for(int ivc = 0; ivc < NVC; ++ivc){
136  wt[ivc+ID1] = bc2 * wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), nei)];
137  wt[ivc+ID2] = bc2 * wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), nei)];
138  wt[ivc+ID3] = bc2 * wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), nei)];
139  wt[ivc+ID4] = bc2 * wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), nei)];
140  }
141  }
142  mult_wilson_ypb(vL, ut, wt);
143  }
144 
145  if ((iy > 0) || (do_comm[idir] == 0)) {
146  int iy2 = (iy - 1 + Ny) % Ny;
147  int nei = ix + Nx * (iy2 + Ny * izt);
148  real_t bc2 = 1.0;
149  if(iy == 0) bc2 = bc[1];
150 
151  real_t ut[NDF];
152  load_u(ut, up, nei + Nst_pad * (1-ieo + 2 * idir));
153 
154  real_t wt[NVCD];
155  if(jgm5 == 0){
156  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
157  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
158  }
159  }else{
160  for(int ivc = 0; ivc < NVC; ++ivc){
161  wt[ivc+ID1] = bc2 * wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), nei)];
162  wt[ivc+ID2] = bc2 * wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), nei)];
163  wt[ivc+ID3] = bc2 * wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), nei)];
164  wt[ivc+ID4] = bc2 * wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), nei)];
165  }
166  }
167  mult_wilson_ymb(vL, ut, wt);
168  }
169 
170  idir = 2;
171 
172  if ((iz < Nz-1) || (do_comm[idir] == 0)) {
173  int iz2 = (iz + 1) % Nz;
174  int nei = ixy + Nxy * (iz2 + Nz * it);
175  real_t bc2 = 1.0;
176  if(iz == Nz-1) bc2 = bc[2];
177 
178  real_t ut[NDF];
179  load_u(ut, up, site + Nst_pad * (ieo + 2 * idir));
180 
181  real_t wt[NVCD];
182  if(jgm5 == 0){
183  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
184  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
185  }
186  }else{
187  for(int ivc = 0; ivc < NVC; ++ivc){
188  wt[ivc+ID1] = bc2 * wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), nei)];
189  wt[ivc+ID2] = bc2 * wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), nei)];
190  wt[ivc+ID3] = bc2 * wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), nei)];
191  wt[ivc+ID4] = bc2 * wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), nei)];
192  }
193  }
194  mult_wilson_zpb(vL, ut, wt);
195  }
196 
197  if ((iz > 0) || (do_comm[idir] == 0)) {
198  int iz2 = (iz - 1 + Nz) % Nz;
199  int nei = ixy + Nxy * (iz2 + Nz * it);
200  real_t bc2 = 1.0;
201  if(iz == 0) bc2 = bc[2];
202 
203  real_t ut[NDF];
204  load_u(ut, up, nei + Nst_pad * (1-ieo + 2 * idir));
205 
206  real_t wt[NVCD];
207  if(jgm5 == 0){
208  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
209  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
210  }
211  }else{
212  for(int ivc = 0; ivc < NVC; ++ivc){
213  wt[ivc+ID1] = bc2 * wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), nei)];
214  wt[ivc+ID2] = bc2 * wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), nei)];
215  wt[ivc+ID3] = bc2 * wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), nei)];
216  wt[ivc+ID4] = bc2 * wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), nei)];
217  }
218  }
219  mult_wilson_zmb(vL, ut, wt);
220  }
221 
222  idir = 3;
223 
224  if ((it < Nt-1) || (do_comm[idir] == 0)) {
225  int it2 = (it + 1) % Nt;
226  int nei = ixyz + Nxyz * it2;
227  real_t bc2 = 1.0;
228  if(it == Nt-1) bc2 = bc[3];
229 
230  real_t ut[NDF];
231  load_u(ut, up, site + Nst_pad * (ieo + 2 * idir));
232 
233  real_t wt[NVCD];
234  if(jgm5 == 0){
235  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
236  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
237  }
238  }else{
239  for(int ivc = 0; ivc < NVC; ++ivc){
240  wt[ivc+ID1] = bc2 * wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), nei)];
241  wt[ivc+ID2] = bc2 * wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), nei)];
242  wt[ivc+ID3] = bc2 * wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), nei)];
243  wt[ivc+ID4] = bc2 * wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), nei)];
244  }
245  }
246  mult_wilson_tpb_dirac(vL, ut, wt);
247  }
248 
249  if ((it > 0) || (do_comm[idir] == 0)) {
250  int it2 = (it - 1 + Nt) % Nt;
251  int nei = ixyz + Nxyz * it2;
252  real_t bc2 = 1.0;
253  if(it == 0) bc2 = bc[3];
254 
255  real_t ut[NDF];
256  load_u(ut, up, nei + Nst_pad * (1-ieo + 2 * idir));
257 
258  real_t wt[NVCD];
259  if(jgm5 == 0){
260  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
261  wt[ivcd] = bc2 * wp[IDX2(Nin5, (ivcd + NVCD * is), nei)];
262  }
263  }else{
264  for(int ivc = 0; ivc < NVC; ++ivc){
265  wt[ivc+ID1] = bc2 * wp[IDX2(Nin5, (ivc+ID3 + NVCD * is), nei)];
266  wt[ivc+ID2] = bc2 * wp[IDX2(Nin5, (ivc+ID4 + NVCD * is), nei)];
267  wt[ivc+ID3] = bc2 * wp[IDX2(Nin5, (ivc+ID1 + NVCD * is), nei)];
268  wt[ivc+ID4] = bc2 * wp[IDX2(Nin5, (ivc+ID2 + NVCD * is), nei)];
269  }
270  }
271  mult_wilson_tmb_dirac(vL, ut, wt);
272  }
273 
274  for(int ivcd = 0; ivcd < NVCD; ++ivcd){
275  vp[IDX2(Nin5, (ivcd + NVCD * is), site)] = vL[ivcd];
276  }
277 
278  } // is loop
279 
280  }
281  } // site loop
282 
283  } // acc parallel
284  } // acc data
285 
286 }
287 
288 #endif
289 //============================================================END=====
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
NVCD
#define NVCD
Definition: define_params_SU3.h:20
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
ID2
#define ID2
Definition: fopr_Wilson_impl_SU2-inc.h:19
ID4
#define ID4
Definition: fopr_Wilson_impl_SU2-inc.h:21
iyzt
int iyzt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:14
ix
int ix
Definition: mult_Wilson_xyz_openacc-inc.h:16
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
mult_domainwall_5din_eo_hopb_dirac_4d
void mult_domainwall_5din_eo_hopb_dirac_4d(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_4d_openacc-inc.h:14
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
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
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
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
iz
int iz
Definition: mult_Wilson_xyz_openacc-inc.h:462
NDIM
#define NDIM
Definition: contract_4spinor.cpp:18
ixyz
int ixyz
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:13