Bridge++  Ver.2.1.3
mult_Clover_openacc-inc.h
Go to the documentation of this file.
1 
9 // This code explicitly assumes SU(3) gauge group.
10 
11 #define MULT_GXr(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_GXi(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 #define MULT_GDXr(u0,u1,u2,u3,u4,u5,v0,v1,v2,v3,v4,v5) (u0*v0+u1*v1 + u2*v2+u3*v3 + u4*v4+u5*v5)
14 #define MULT_GDXi(u0,u1,u2,u3,u4,u5,v0,v1,v2,v3,v4,v5) (u0*v1-u1*v0 + u2*v3-u3*v2 + u4*v5-u5*v4)
15 
16 #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)
17 #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)
18 
19 #define EXT_IMG_R(v1r,v1i,v2r,v2i,w1r,w1i,w2r,w2i) (v1r*w2r - v1i*w2i - v2r*w1r + v2i*w1i)
20 #define EXT_IMG_I(v1r,v1i,v2r,v2i,w1r,w1i,w2r,w2i) (- v1r*w2i - v1i*w2r + v2r*w1i + v2i*w1r)
21 
22 
23 //====================================================================
24 void mult_clover_D_dirac(real_t *RESTRICT v2, real_t *RESTRICT u,
25  real_t *RESTRICT ct, real_t *RESTRICT v1,
26  real_t kappa, int *Nsize, int *bc, int flag)
27 { // flag=0: D, flag=1: H
28 
29  int Nx = Nsize[0];
30  int Ny = Nsize[1];
31  int Nz = Nsize[2];
32  int Nt = Nsize[3];
33  int Nst = Nx * Ny * Nz * Nt;
34  int Nst_pad = CEIL_NWP(Nst);
35 
36  int nvst = NVC * ND * Nst_pad;
37  int ngst = NDF * Nst_pad * 4;
38  int ncst = NDF * Nst_pad * 8;
39 
40  real_t *RESTRICT u_up = u;
41  real_t *RESTRICT u_dn = u;
42 
43 #pragma acc data present(v1[0:nvst], u_up[0:ngst], u_dn[0:ngst], \
44  ct[0:ncst], v2[0:nvst]) \
45  copyin(bc[0:4], kappa, Nx, Ny, Nz, Nt, Nst, Nst_pad)
46 
47 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
48  {
49 
50 #pragma acc loop gang worker vector
51  for(int site = 0; site < Nst; ++site){
52  int Nxy = Nx * Ny;
53  int Nxyz = Nx * Ny * Nz;
54 
55  real_t bc2;
56 
57  real_t u_0, u_1, u_2, u_3, u_4, u_5;
58  real_t u_6, u_7, u_8, u_9, u10, u11;
59  real_t u12, u13, u14, u15, u16, u17;
60  real_t vt1_0, vt1_1, vt1_2, vt1_3, vt1_4, vt1_5;
61  real_t vt2_0, vt2_1, vt2_2, vt2_3, vt2_4, vt2_5;
62  real_t wt1r, wt1i, wt2r, wt2i;
63 
64  real_t v2_01, v2_11, v2_21, v2_31, v2_41, v2_51;
65  real_t v2_02, v2_12, v2_22, v2_32, v2_42, v2_52;
66  real_t v2_03, v2_13, v2_23, v2_33, v2_43, v2_53;
67  real_t v2_04, v2_14, v2_24, v2_34, v2_44, v2_54;
68 
69 #include "inc/mult_Wilson_xyz_openacc-inc.h"
70 
71 #include "inc/mult_Wilson_t_dirac_openacc-inc.h"
72 
73 #include "inc/mult_Clover_csw_dirac_openacc-inc.h"
74 
75  // write back to global memory
76  if(flag == 0){ // D
77  v2[IDX2_SP_R(0, 0, site)] = v2_01;
78  v2[IDX2_SP_I(0, 0, site)] = v2_11;
79  v2[IDX2_SP_R(1, 0, site)] = v2_21;
80  v2[IDX2_SP_I(1, 0, site)] = v2_31;
81  v2[IDX2_SP_R(2, 0, site)] = v2_41;
82  v2[IDX2_SP_I(2, 0, site)] = v2_51;
83 
84  v2[IDX2_SP_R(0, 1, site)] = v2_02;
85  v2[IDX2_SP_I(0, 1, site)] = v2_12;
86  v2[IDX2_SP_R(1, 1, site)] = v2_22;
87  v2[IDX2_SP_I(1, 1, site)] = v2_32;
88  v2[IDX2_SP_R(2, 1, site)] = v2_42;
89  v2[IDX2_SP_I(2, 1, site)] = v2_52;
90 
91  v2[IDX2_SP_R(0, 2, site)] = v2_03;
92  v2[IDX2_SP_I(0, 2, site)] = v2_13;
93  v2[IDX2_SP_R(1, 2, site)] = v2_23;
94  v2[IDX2_SP_I(1, 2, site)] = v2_33;
95  v2[IDX2_SP_R(2, 2, site)] = v2_43;
96  v2[IDX2_SP_I(2, 2, site)] = v2_53;
97 
98  v2[IDX2_SP_R(0, 3, site)] = v2_04;
99  v2[IDX2_SP_I(0, 3, site)] = v2_14;
100  v2[IDX2_SP_R(1, 3, site)] = v2_24;
101  v2[IDX2_SP_I(1, 3, site)] = v2_34;
102  v2[IDX2_SP_R(2, 3, site)] = v2_44;
103  v2[IDX2_SP_I(2, 3, site)] = v2_54;
104  }else{ // H
105  v2[IDX2_SP_R(0, 0, site)] = v2_03;
106  v2[IDX2_SP_I(0, 0, site)] = v2_13;
107  v2[IDX2_SP_R(1, 0, site)] = v2_23;
108  v2[IDX2_SP_I(1, 0, site)] = v2_33;
109  v2[IDX2_SP_R(2, 0, site)] = v2_43;
110  v2[IDX2_SP_I(2, 0, site)] = v2_53;
111 
112  v2[IDX2_SP_R(0, 1, site)] = v2_04;
113  v2[IDX2_SP_I(0, 1, site)] = v2_14;
114  v2[IDX2_SP_R(1, 1, site)] = v2_24;
115  v2[IDX2_SP_I(1, 1, site)] = v2_34;
116  v2[IDX2_SP_R(2, 1, site)] = v2_44;
117  v2[IDX2_SP_I(2, 1, site)] = v2_54;
118 
119  v2[IDX2_SP_R(0, 2, site)] = v2_01;
120  v2[IDX2_SP_I(0, 2, site)] = v2_11;
121  v2[IDX2_SP_R(1, 2, site)] = v2_21;
122  v2[IDX2_SP_I(1, 2, site)] = v2_31;
123  v2[IDX2_SP_R(2, 2, site)] = v2_41;
124  v2[IDX2_SP_I(2, 2, site)] = v2_51;
125 
126  v2[IDX2_SP_R(0, 3, site)] = v2_02;
127  v2[IDX2_SP_I(0, 3, site)] = v2_12;
128  v2[IDX2_SP_R(1, 3, site)] = v2_22;
129  v2[IDX2_SP_I(1, 3, site)] = v2_32;
130  v2[IDX2_SP_R(2, 3, site)] = v2_42;
131  v2[IDX2_SP_I(2, 3, site)] = v2_52;
132  }
133 
134  }
135 
136  }
137 
138 }
139 
140 //====================================================================
141 void mult_clover_D_chiral(real_t *RESTRICT v2, real_t *RESTRICT u,
142  real_t *RESTRICT ct, real_t *RESTRICT v1,
143  real_t kappa, int *Nsize, int *bc, int flag)
144 { // flag=0: D, flag=1: H
145 
146  int Nx = Nsize[0];
147  int Ny = Nsize[1];
148  int Nz = Nsize[2];
149  int Nt = Nsize[3];
150  int Nst = Nx * Ny * Nz * Nt;
151  int Nst_pad = CEIL_NWP(Nst);
152 
153  int nvst = NVC * ND * Nst_pad;
154  int ngst = NDF * Nst_pad * 4;
155  int ncst = NDF * Nst_pad * 8;
156 
157  real_t *RESTRICT u_up = u;
158  real_t *RESTRICT u_dn = u;
159 
160 #pragma acc data present(v1[0:nvst], u_up[0:ngst], u_dn[0:ngst], \
161  ct[0:ncst], v2[0:nvst]) \
162  copyin(bc[0:4], kappa, Nx, Ny, Nz, Nt,Nst, Nst_pad)
163 
164 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
165  {
166 
167 #pragma acc loop gang worker vector
168  for(int site = 0; site < Nst; ++site){
169  int Nxy = Nx * Ny;
170  int Nxyz = Nx * Ny * Nz;
171 
172  real_t bc2;
173 
174  real_t u_0, u_1, u_2, u_3, u_4, u_5;
175  real_t u_6, u_7, u_8, u_9, u10, u11;
176  real_t u12, u13, u14, u15, u16, u17;
177  real_t vt1_0, vt1_1, vt1_2, vt1_3, vt1_4, vt1_5;
178  real_t vt2_0, vt2_1, vt2_2, vt2_3, vt2_4, vt2_5;
179  real_t wt1r, wt1i, wt2r, wt2i;
180 
181  real_t v2_01, v2_11, v2_21, v2_31, v2_41, v2_51;
182  real_t v2_02, v2_12, v2_22, v2_32, v2_42, v2_52;
183  real_t v2_03, v2_13, v2_23, v2_33, v2_43, v2_53;
184  real_t v2_04, v2_14, v2_24, v2_34, v2_44, v2_54;
185 
186 #include "inc/mult_Wilson_xyz_openacc-inc.h"
187 
188 #include "inc/mult_Wilson_t_chiral_openacc-inc.h"
189 
190 #include "inc/mult_Clover_csw_chiral_openacc-inc.h"
191 
192  // aypx and write back to global memory
193  if(flag == 0){ // D
194  v2[IDX2_SP_R(0, 0, site)] = v2_01;
195  v2[IDX2_SP_I(0, 0, site)] = v2_11;
196  v2[IDX2_SP_R(1, 0, site)] = v2_21;
197  v2[IDX2_SP_I(1, 0, site)] = v2_31;
198  v2[IDX2_SP_R(2, 0, site)] = v2_41;
199  v2[IDX2_SP_I(2, 0, site)] = v2_51;
200 
201  v2[IDX2_SP_R(0, 1, site)] = v2_02;
202  v2[IDX2_SP_I(0, 1, site)] = v2_12;
203  v2[IDX2_SP_R(1, 1, site)] = v2_22;
204  v2[IDX2_SP_I(1, 1, site)] = v2_32;
205  v2[IDX2_SP_R(2, 1, site)] = v2_42;
206  v2[IDX2_SP_I(2, 1, site)] = v2_52;
207 
208  v2[IDX2_SP_R(0, 2, site)] = v2_03;
209  v2[IDX2_SP_I(0, 2, site)] = v2_13;
210  v2[IDX2_SP_R(1, 2, site)] = v2_23;
211  v2[IDX2_SP_I(1, 2, site)] = v2_33;
212  v2[IDX2_SP_R(2, 2, site)] = v2_43;
213  v2[IDX2_SP_I(2, 2, site)] = v2_53;
214 
215  v2[IDX2_SP_R(0, 3, site)] = v2_04;
216  v2[IDX2_SP_I(0, 3, site)] = v2_14;
217  v2[IDX2_SP_R(1, 3, site)] = v2_24;
218  v2[IDX2_SP_I(1, 3, site)] = v2_34;
219  v2[IDX2_SP_R(2, 3, site)] = v2_44;
220  v2[IDX2_SP_I(2, 3, site)] = v2_54;
221  }else{
222  v2[IDX2_SP_R(0, 0, site)] = v2_01;
223  v2[IDX2_SP_I(0, 0, site)] = v2_11;
224  v2[IDX2_SP_R(1, 0, site)] = v2_21;
225  v2[IDX2_SP_I(1, 0, site)] = v2_31;
226  v2[IDX2_SP_R(2, 0, site)] = v2_41;
227  v2[IDX2_SP_I(2, 0, site)] = v2_51;
228 
229  v2[IDX2_SP_R(0, 1, site)] = v2_02;
230  v2[IDX2_SP_I(0, 1, site)] = v2_12;
231  v2[IDX2_SP_R(1, 1, site)] = v2_22;
232  v2[IDX2_SP_I(1, 1, site)] = v2_32;
233  v2[IDX2_SP_R(2, 1, site)] = v2_42;
234  v2[IDX2_SP_I(2, 1, site)] = v2_52;
235 
236  v2[IDX2_SP_R(0, 2, site)] = v2_03;
237  v2[IDX2_SP_I(0, 2, site)] = v2_13;
238  v2[IDX2_SP_R(1, 2, site)] = v2_23;
239  v2[IDX2_SP_I(1, 2, site)] = v2_33;
240  v2[IDX2_SP_R(2, 2, site)] = v2_43;
241  v2[IDX2_SP_I(2, 2, site)] = v2_53;
242 
243  v2[IDX2_SP_R(0, 3, site)] = v2_04;
244  v2[IDX2_SP_I(0, 3, site)] = v2_14;
245  v2[IDX2_SP_R(1, 3, site)] = v2_24;
246  v2[IDX2_SP_I(1, 3, site)] = v2_34;
247  v2[IDX2_SP_R(2, 3, site)] = v2_44;
248  v2[IDX2_SP_I(2, 3, site)] = v2_54;
249  }
250  }
251 
252  }
253 
254 }
255 
256 //============================================================END=====
u14
u14
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:100
v2_24
v2_24
Definition: mult_Clover_csw_chiral_openacc-inc.h:105
v2_22
v2_22
Definition: mult_Clover_csw_chiral_openacc-inc.h:91
u13
u13
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:99
v2_53
v2_53
Definition: mult_Clover_csw_chiral_openacc-inc.h:101
v2_01
v2_01
Definition: mult_Clover_csw_chiral_openacc-inc.h:82
u11
u11
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:74
v2_52
v2_52
Definition: mult_Clover_csw_chiral_openacc-inc.h:94
v2_14
v2_14
Definition: mult_Clover_csw_chiral_openacc-inc.h:104
vt2_1
vt2_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:40
v2_04
v2_04
Definition: mult_Clover_csw_chiral_openacc-inc.h:103
u_0
u_0
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:46
v2_11
v2_11
Definition: mult_Clover_csw_chiral_openacc-inc.h:83
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
vt2_5
vt2_5
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:44
u_9
u_9
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:72
vt1_2
vt1_2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:34
mult_Clover_csw_dirac_openacc-inc.h
u_5
u_5
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:51
wt1i
wt1i
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:55
v2_13
v2_13
Definition: mult_Clover_csw_chiral_openacc-inc.h:97
vt2_2
vt2_2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:41
wt2r
wt2r
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:57
mult_Wilson_t_dirac_openacc-inc.h
v2_23
v2_23
Definition: mult_Clover_csw_chiral_openacc-inc.h:98
u_3
u_3
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:49
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
v2_21
v2_21
Definition: mult_Clover_csw_chiral_openacc-inc.h:84
u_8
u_8
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:71
vt1_0
vt1_0
Definition: mult_Wilson_eo_t_chiral_openacc-inc.h:18
u10
u10
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:73
bc2
int bc2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:62
mult_Clover_csw_chiral_openacc-inc.h
wt1r
wt1r
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:53
wt2i
wt2i
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:59
mult_Wilson_t_chiral_openacc-inc.h
v2_31
v2_31
Definition: mult_Clover_csw_chiral_openacc-inc.h:85
u_2
u_2
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:48
u16
u16
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:102
u_1
u_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:47
u_4
u_4
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:50
v2_32
v2_32
Definition: mult_Clover_csw_chiral_openacc-inc.h:92
ND
#define ND
Definition: field_F_imp_SU2-inc.h:18
u17
u17
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:103
v2_42
v2_42
Definition: mult_Clover_csw_chiral_openacc-inc.h:93
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
v2_03
v2_03
Definition: mult_Clover_csw_chiral_openacc-inc.h:96
mult_clover_D_chiral
void mult_clover_D_chiral(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT ct, real_t *RESTRICT v1, real_t kappa, int *Nsize, int *bc, int flag)
Definition: mult_Clover_openacc-inc.h:141
v2_51
v2_51
Definition: mult_Clover_csw_chiral_openacc-inc.h:87
vt1_5
vt1_5
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:37
v2_12
v2_12
Definition: mult_Clover_csw_chiral_openacc-inc.h:90
u12
u12
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:98
v2_02
v2_02
Definition: mult_Clover_csw_chiral_openacc-inc.h:89
vt2_3
vt2_3
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:42
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
mult_clover_D_dirac
void mult_clover_D_dirac(real_t *RESTRICT v2, real_t *RESTRICT u, real_t *RESTRICT ct, real_t *RESTRICT v1, real_t kappa, int *Nsize, int *bc, int flag)
Definition: mult_Clover_openacc-inc.h:24
v2_34
v2_34
Definition: mult_Clover_csw_chiral_openacc-inc.h:106
v2_41
v2_41
Definition: mult_Clover_csw_chiral_openacc-inc.h:86
vt1_3
vt1_3
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:35
v2_33
v2_33
Definition: mult_Clover_csw_chiral_openacc-inc.h:99
v2_44
v2_44
Definition: mult_Clover_csw_chiral_openacc-inc.h:107
IDX2_SP_R
#define IDX2_SP_R(ic, id, ist)
Definition: define_index.h:31
u15
u15
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:101
vt2_4
vt2_4
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:43
u_6
u_6
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:69
v2_43
v2_43
Definition: mult_Clover_csw_chiral_openacc-inc.h:100
vt2_0
vt2_0
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:39
IDX2_SP_I
#define IDX2_SP_I(ic, id, ist)
Definition: define_index.h:32
mult_Wilson_xyz_openacc-inc.h
u_7
u_7
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:70
v2_54
v2_54
Definition: mult_Clover_csw_chiral_openacc-inc.h:108
vt1_4
vt1_4
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:36
vt1_1
vt1_1
Definition: mult_Domainwall_eo_t_dirac_openacc-inc.h:33