Bridge++  Ver.2.1.3
mult_Wilson_1xyz_openacc-inc.h
Go to the documentation of this file.
1 
9  // int idir = 0;
10  if(do_comm[0] > 0){
11 
12 #pragma acc parallel async \
13  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
14  {
15  int Nyzt = Ny * Nz * Nt;
16 
17 #pragma acc loop
18  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
19  int ix = 0;
20  int ist = ix + Nx * iyzt;
21  real_t bc2 = bc[0];
22  for(int ic = 0; ic < NC; ++ic){
23  real_t vt1[2], vt2[2];
24  vt1[0] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
25  vt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
26  vt2[0] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
27  vt2[1] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
28  buf_xp[IDXBF_R(ic, 0, iyzt)] = bc2 * vt1[0];
29  buf_xp[IDXBF_I(ic, 0, iyzt)] = bc2 * vt1[1];
30  buf_xp[IDXBF_R(ic, 1, iyzt)] = bc2 * vt2[0];
31  buf_xp[IDXBF_I(ic, 1, iyzt)] = bc2 * vt2[1];
32  }
33  }
34 
35  }
36 
37 #pragma acc parallel async \
38  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
39  {
40  int Nyzt = Ny * Nz * Nt;
41 
42 #pragma acc loop
43  for(int iyzt = 0; iyzt < Nyzt; ++iyzt){
44  int ix = Nx-1;
45  int ist = ix + Nx * iyzt;
46 
47  real_t vt1[NVC], vt2[NVC], ut[NDF];
48  for(int ic = 0; ic < NC; ++ic){
49  int icr = 2*ic;
50  int ici = 2*ic + 1;
51  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
52  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
53  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
54  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
55  }
56 
57  real_t wt1[2], wt2[2];
58 
59  for(int ic = 0; ic < NC; ++ic){
60 
61  for(int ic2 = 0; ic2 < NC; ++ic2){
62  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, ist)];
63  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, ist)];
64  }
65 
66  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
67  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
68  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
69  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
70  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
71  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
72  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
73  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
74 
75  buf_xm[IDXBF_R(ic, 0, iyzt)] = wt1[0];
76  buf_xm[IDXBF_I(ic, 0, iyzt)] = wt1[1];
77  buf_xm[IDXBF_R(ic, 1, iyzt)] = wt2[0];
78  buf_xm[IDXBF_I(ic, 1, iyzt)] = wt2[1];
79 
80  }
81 
82  }
83 
84  }
85 
86  } // do_comm[0]
87 
88  // idir = 1;
89  if(do_comm[1] > 0){
90 
91 #pragma acc parallel async \
92  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
93  {
94  int Nzt = Nz * Nt;
95 
96 #pragma acc loop collapse(2)
97  for(int izt = 0; izt < Nzt; ++izt){
98  for(int ix = 0; ix < Nx; ++ix){
99  int iy = 0;
100  int ist = ix + Nx * (iy + Ny * izt);
101  int ixzt = ix + Nx * izt;
102  real_t bc2 = bc[1];
103 
104  for(int ic = 0; ic < NC; ++ic){
105  real_t vt1[2], vt2[2];
106  vt1[0] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
107  vt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
108  vt2[0] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
109  vt2[1] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
110  buf_yp[IDXBF_R(ic, 0, ixzt)] = bc2 * vt1[0];
111  buf_yp[IDXBF_I(ic, 0, ixzt)] = bc2 * vt1[1];
112  buf_yp[IDXBF_R(ic, 1, ixzt)] = bc2 * vt2[0];
113  buf_yp[IDXBF_I(ic, 1, ixzt)] = bc2 * vt2[1];
114  }
115 
116  }
117  }
118 
119  }
120 
121 #pragma acc parallel async \
122  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
123  {
124  int Nzt = Nz * Nt;
125 
126 #pragma acc loop collapse(2)
127  for(int izt = 0; izt < Nzt; ++izt){
128  for(int ix = 0; ix < Nx; ++ix){
129  int iy = Ny-1;
130  int ist = ix + Nx * (iy + Ny*izt);
131  int istu = ist + Nst_pad;
132  int ixzt = ix + Nx * izt;
133 
134  real_t vt1[NVC], vt2[NVC];
135  for(int ic = 0; ic < NC; ++ic){
136  int icr = 2*ic;
137  int ici = 2*ic + 1;
138  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
139  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
140  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
141  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
142  }
143 
144  for(int ic = 0; ic < NC; ++ic){
145 
146  real_t ut[NVC];
147  for(int ic2 = 0; ic2 < NC; ++ic2){
148  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
149  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
150  }
151 
152  real_t wt1[2], wt2[2];
153  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
154  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
155  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
156  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
157  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
158  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
159  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
160  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
161 
162  buf_ym[IDXBF_R(ic, 0, ixzt)] = wt1[0];
163  buf_ym[IDXBF_I(ic, 0, ixzt)] = wt1[1];
164  buf_ym[IDXBF_R(ic, 1, ixzt)] = wt2[0];
165  buf_ym[IDXBF_I(ic, 1, ixzt)] = wt2[1];
166 
167  }
168 
169  }
170  }
171 
172  }
173 
174  } // do_comm[1]
175 
176  // idir = 2;
177  if(do_comm[2] > 0){
178 
179 #pragma acc parallel async \
180  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
181  {
182  int Nxy = Nx * Ny;
183 
184 #pragma acc loop collapse(2)
185  for(int it = 0; it < Nt; ++it){
186  for(int ixy = 0; ixy < Nxy; ++ixy){
187  int iz = 0;
188  int ist = ixy + Nxy * (iz + Nz * it);
189  int ixyt = ixy + Nxy * it;
190  real_t bc2 = bc[2];
191 
192  for(int ic = 0; ic < NC; ++ic){
193  real_t wt1[2], wt2[2];
194  wt1[0] = v1[IDX2_SP_R(ic, 0, ist)] - v1[IDX2_SP_I(ic, 2, ist)];
195  wt1[1] = v1[IDX2_SP_I(ic, 0, ist)] + v1[IDX2_SP_R(ic, 2, ist)];
196  wt2[0] = v1[IDX2_SP_R(ic, 1, ist)] + v1[IDX2_SP_I(ic, 3, ist)];
197  wt2[1] = v1[IDX2_SP_I(ic, 1, ist)] - v1[IDX2_SP_R(ic, 3, ist)];
198  buf_zp[IDXBF_R(ic, 0, ixyt)] = bc2 * wt1[0];
199  buf_zp[IDXBF_I(ic, 0, ixyt)] = bc2 * wt1[1];
200  buf_zp[IDXBF_R(ic, 1, ixyt)] = bc2 * wt2[0];
201  buf_zp[IDXBF_I(ic, 1, ixyt)] = bc2 * wt2[1];
202  }
203 
204  }
205  }
206 
207  }
208 
209 #pragma acc parallel async \
210  num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
211  {
212  int Nxy = Nx * Ny;
213 
214 #pragma acc loop collapse(2)
215  for(int it = 0; it < Nt; ++it){
216  for(int ixy = 0; ixy < Nxy; ++ixy){
217  int iz = Nz-1;
218  int ist = ixy + Nxy * (iz + Nz * it);
219  int istu = ist + Nst_pad * 2;
220  int ixyt = ixy + Nxy * it;
221 
222  real_t vt1[NVC], vt2[NVC];
223  for(int ic = 0; ic < NC; ++ic){
224  int icr = 2*ic;
225  int ici = 2*ic + 1;
226  vt1[icr] = v1[IDX2_SP_R(ic, 0, ist)] + v1[IDX2_SP_I(ic, 2, ist)];
227  vt1[ici] = v1[IDX2_SP_I(ic, 0, ist)] - v1[IDX2_SP_R(ic, 2, ist)];
228  vt2[icr] = v1[IDX2_SP_R(ic, 1, ist)] - v1[IDX2_SP_I(ic, 3, ist)];
229  vt2[ici] = v1[IDX2_SP_I(ic, 1, ist)] + v1[IDX2_SP_R(ic, 3, ist)];
230  }
231 
232  for(int ic = 0; ic < NC; ++ic){
233 
234  real_t ut[NVC];
235  for(int ic2 = 0; ic2 < NC; ++ic2){
236  ut[2*ic2 ] = u[IDX2_G_R(ic, ic2, istu)];
237  ut[2*ic2+1] = - u[IDX2_G_I(ic, ic2, istu)];
238  }
239 
240  real_t wt1[2], wt2[2];
241  wt1[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
242  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
243  wt1[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
244  vt1[0],vt1[1],vt1[2],vt1[3],vt1[4],vt1[5]);
245  wt2[0] = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
246  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
247  wt2[1] = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
248  vt2[0],vt2[1],vt2[2],vt2[3],vt2[4],vt2[5]);
249  buf_zm[IDXBF_R(ic, 0, ixyt)] = wt1[0];
250  buf_zm[IDXBF_I(ic, 0, ixyt)] = wt1[1];
251  buf_zm[IDXBF_R(ic, 1, ixyt)] = wt2[0];
252  buf_zm[IDXBF_I(ic, 1, ixyt)] = wt2[1];
253 
254  }
255 
256  }
257  }
258 
259  }
260 
261  } // do_comm[2]
262 
263 //============================================================END=====
iy
int iy
Definition: mult_Wilson_xyz_openacc-inc.h:239
ixy
int ixy
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:513
izt
int izt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:262
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
IDXBF_I
#define IDXBF_I(ic, id, ist)
Definition: define_index.h:42
IDX2_G_I
#define IDX2_G_I(ic1, ic2, ist)
Definition: define_index.h:52
iyzt
int iyzt
Definition: mult_Domainwall_eo_xyz_openacc-inc.h:14
ix
int ix
Definition: mult_Wilson_xyz_openacc-inc.h:16
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
it
int it
Definition: mult_Wilson_xyz_openacc-inc.h:461
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
MULT_UV_I
#define MULT_UV_I(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_Wilson_inline_openacc-inc.h:14
IDXBF_R
#define IDXBF_R(ic, id, ist)
Definition: define_index.h:41
MULT_UV_R
#define MULT_UV_R(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_Wilson_inline_openacc-inc.h:13
IDX2_SP_R
#define IDX2_SP_R(ic, id, ist)
Definition: define_index.h:31
iz
int iz
Definition: mult_Wilson_xyz_openacc-inc.h:462
IDX2_G_R
#define IDX2_G_R(ic1, ic2, ist)
Definition: define_index.h:51
vt2
real_t vt2
Definition: mult_Staggered_uvdn1_openacc-inc.h:8
IDX2_SP_I
#define IDX2_SP_I(ic, id, ist)
Definition: define_index.h:32