Bridge++  Ver.2.1.3
mult_CloverTerm_openacc-inc.h
Go to the documentation of this file.
1 
10 #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)
11 #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)
12 
13 
14 //====================================================================
15 void mult_csw_dirac(real_t *RESTRICT v2, real_t *RESTRICT u,
16  real_t *RESTRICT v1, int *Nsize, int iflag)
17  // iflag = 0: mult, iflag = 1: multadd
18 {
19  int Nst = Nsize[0] * Nsize[1] * Nsize[2] * Nsize[3];
20  int Nst_pad = CEIL_NWP(Nst);
21 
22  int size = NVC * ND * Nst_pad;
23  int size_u = NDF * Nst_pad * ND * ND2;
24 
25 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]) \
26  copyin(Nst, Nst_pad, iflag)
27 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
28  {
29 
30 #pragma acc loop gang worker vector
31  for(int ist = 0; ist < Nst; ++ist){
32 
33  real_t ut[NVC], vt[NVC * ND], wt[NVC * ND];
34 
35  if(iflag == 0){
36  for(int id = 0; id < ND; ++id){
37  for(int ivc = 0; ivc < NVC; ++ivc){
38  wt[ivc + NVC*id] = 0.0;
39  }
40  }
41  }else{
42  for(int id = 0; id < ND; ++id){
43  for(int ic = 0; ic < NC; ++ic){
44  wt[2*ic + NVC*id] = v2[IDX2_SP_R(ic, id, ist)];
45  wt[2*ic+1 + NVC*id] = v2[IDX2_SP_I(ic, id, ist)];
46  }
47  }
48  }
49 
50  for(int id = 0; id < ND; ++id){
51  for(int ic = 0; ic < NC; ++ic){
52  vt[2*ic + NVC*id] = v1[IDX2_SP_R(ic, id, ist)];
53  vt[2*ic+1 + NVC*id] = v1[IDX2_SP_I(ic, id, ist)];
54  }
55  }
56 
57  for(int jd = 0; jd < ND2; ++jd){
58  for(int id = 0; id < ND; ++id){
59  int igst = ist + Nst_pad * (id + ND * jd);
60 
61  for(int ic2 = 0; ic2 < NC; ++ic2){
62 
63  for(int ic1 = 0; ic1 < NC; ++ic1){
64  ut[2*ic1 ] = u[IDX2_G_R(ic1, ic2, igst)];
65  ut[2*ic1+1] = u[IDX2_G_I(ic1, ic2, igst)];
66  }
67 
68  int id2 = (id + ND2) % ND;
69  int j = NVC * id;
70  int k = NVC * id2;
71  real_t wt1r, wt1i, wt2r, wt2i;
72  wt1r = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
73  vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j]);
74 
75  wt1i = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
76  vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j]);
77 
78  wt2r = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
79  vt[0+k], vt[1+k], vt[2+k], vt[3+k], vt[4+k], vt[5+k]);
80 
81  wt2i = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
82  vt[0+k], vt[1+k], vt[2+k], vt[3+k], vt[4+k], vt[5+k]);
83 
84  int jd2 = jd + ND2;
85  wt[2*ic2 + NVC * jd] += wt1r;
86  wt[2*ic2+1 + NVC * jd] += wt1i;
87  wt[2*ic2 + NVC * jd2] += wt2r;
88  wt[2*ic2+1 + NVC * jd2] += wt2i;
89 
90  }
91  }
92  }
93 
94  for(int id = 0; id < ND; ++id){
95  for(int ic = 0; ic < NC; ++ic){
96  v2[IDX2_SP_R(ic, id, ist)] = wt[2*ic + NVC * id];
97  v2[IDX2_SP_I(ic, id, ist)] = wt[2*ic+1 + NVC * id];
98  }
99  }
100 
101  }
102 
103  } // acc parallel
104 
105  //#pragma omp barrier
106 
107 }
108 
109 //====================================================================
110 void mult_csw_chiral(real_t *RESTRICT v2, real_t *RESTRICT u,
111  real_t *RESTRICT v1, int *Nsize, int iflag)
112  // iflag = 0: mult, iflag = 1: multadd
113 {
114  int Nst = Nsize[0] * Nsize[1] * Nsize[2] * Nsize[3];
115  int Nst_pad = CEIL_NWP(Nst);
116 
117  int size = NVC * ND * Nst_pad;
118  int size_u = NDF * Nst_pad * ND * ND2;
119 
120 #pragma acc data present(v2[0:size], v1[0:size], u[0:size_u]) \
121  copyin(Nst, Nst_pad, iflag)
122 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
123  {
124 
125 #pragma acc loop gang worker vector
126  for(int ist = 0; ist < Nst; ++ist){
127 
128  real_t ut[NVC], vt[NVC * ND], wt[NVC * ND];
129 
130  if(iflag == 0){
131  for(int id = 0; id < ND; ++id){
132  for(int ivc = 0; ivc < NVC; ++ivc){
133  wt[ivc + NVC*id] = 0.0;
134  }
135  }
136  }else{
137  for(int id = 0; id < ND; ++id){
138  for(int ic = 0; ic < NC; ++ic){
139  wt[2*ic + NVC*id] = v2[IDX2_SP_R(ic, id, ist)];
140  wt[2*ic+1 + NVC*id] = v2[IDX2_SP_I(ic, id, ist)];
141  }
142  }
143  }
144 
145  for(int id = 0; id < ND; ++id){
146  for(int ic = 0; ic < NC; ++ic){
147  vt[2*ic + NVC*id] = v1[IDX2_SP_R(ic, id, ist)];
148  vt[2*ic+1 + NVC*id] = v1[IDX2_SP_I(ic, id, ist)];
149  }
150  }
151 
152  for(int jd = 0; jd < ND2; ++jd){
153  for(int id = 0; id < ND2; ++id){
154  int igst = ist + Nst_pad * (id + ND2 * jd);
155 
156  for(int ic2 = 0; ic2 < NC; ++ic2){
157 
158  for(int ic1 = 0; ic1 < NC; ++ic1){
159  ut[2*ic1 ] = u[IDX2_G_R(ic1, ic2, igst)];
160  ut[2*ic1+1] = u[IDX2_G_I(ic1, ic2, igst)];
161  }
162 
163  int j = NVC * id;
164  real_t wt1r, wt1i;
165  wt1r = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
166  vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j]);
167 
168  wt1i = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
169  vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j]);
170 
171  wt[2*ic2 + NVC*jd] += wt1r;
172  wt[2*ic2+1 + NVC*jd] += wt1i;
173 
174  }
175 
176  }
177  }
178 
179  for(int jd = 0; jd < ND2; ++jd){
180  for(int id = 0; id < ND2; ++id){
181  int igst = ist + Nst_pad * (id + ND2*jd + ND);
182 
183  for(int ic2 = 0; ic2 < NC; ++ic2){
184 
185  for(int ic1 = 0; ic1 < NC; ++ic1){
186  ut[2*ic1 ] = u[IDX2_G_R(ic1, ic2, igst)];
187  ut[2*ic1+1] = u[IDX2_G_I(ic1, ic2, igst)];
188  }
189 
190  int j = NVC * (id + ND2);
191  real_t wt1r, wt1i;
192  wt1r = MULT_UV_R(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
193  vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j]);
194 
195  wt1i = MULT_UV_I(ut[0], ut[1], ut[2], ut[3], ut[4], ut[5],
196  vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j]);
197 
198  int jd2 = jd + ND2;
199  wt[2*ic2 + NVC * jd2] += wt1r;
200  wt[2*ic2+1 + NVC * jd2] += wt1i;
201 
202  }
203 
204  }
205  }
206 
207  for(int id = 0; id < ND; ++id){
208  for(int ic = 0; ic < NC; ++ic){
209  v2[IDX2_SP_R(ic, id, ist)] = wt[2*ic + NVC * id];
210  v2[IDX2_SP_I(ic, id, ist)] = wt[2*ic+1 + NVC * id];
211  }
212  }
213 
214  }
215 
216  } // acc parallel
217 
218  //#pragma omp barrier
219 
220 }
221 
222 //============================================================END=====
wt
real_t wt[NVCD]
Definition: mult_Clover_csw_chiral_openacc-inc.h:9
mult_csw_dirac
void mult_csw_dirac(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int iflag)
Definition: mult_CloverTerm_openacc-inc.h:15
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
wt1i
wt1i
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:55
wt2r
wt2r
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:57
mult_csw_chiral
void mult_csw_chiral(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT v1, int *Nsize, int iflag)
Definition: mult_CloverTerm_openacc-inc.h:110
MULT_UV_R
#define MULT_UV_R(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_CloverTerm_openacc-inc.h:10
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
IDX2_G_I
#define IDX2_G_I(ic1, ic2, ist)
Definition: define_index.h:52
MULT_UV_I
#define MULT_UV_I(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: mult_CloverTerm_openacc-inc.h:11
NC
#define NC
Definition: field_F_imp_SU2-inc.h:15
wt1r
wt1r
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:53
wt2i
wt2i
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:59
ND
#define ND
Definition: field_F_imp_SU2-inc.h:18
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
IDX2_SP_R
#define IDX2_SP_R(ic, id, ist)
Definition: define_index.h:31
IDX2_G_R
#define IDX2_G_R(ic1, ic2, ist)
Definition: define_index.h:51
ND2
#define ND2
Definition: define_params_SU3.h:18
IDX2_SP_I
#define IDX2_SP_I(ic, id, ist)
Definition: define_index.h:32