Bridge++  Ver.2.1.3
afield_Gauge_openacc-inc.h
Go to the documentation of this file.
1 
10 #ifndef AFIELD_GAUGE_OPENACC_INC_INCLUDED
11 #define AFIELD_GAUGE_OPENACC_INC_INCLUDED
12 
13 #include "inline/define_params.h"
14 #include "inline/define_index.h"
15 
16 #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)
17 #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)
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 #define MULT_R(v1r, v1i, v2r, v2i) (v1r * v2r - v1i * v2i)
23 #define MULT_I(v1r, v1i, v2r, v2i) (v1r * v2i + v1i * v2r)
24 
25 //====================================================================
26 inline real_t mult_r(const real_t v1r, const real_t v1i,
27  const real_t v2r, const real_t v2i){
28  return v1r * v2r - v1i * v2i;
29 }
30 
31 inline real_t mult_i(const real_t v1r, const real_t v1i,
32  const real_t v2r, const real_t v2i){
33  return v1r * v2i + v1i * v2r;
34 }
35 
36 //====================================================================
37 void multadd_Gnn(real_t *RESTRICT u, const int exu,
38  real_t *RESTRICT v, const int exv,
39  real_t *RESTRICT w, const int exw,
40  const real_t a, const int Nst)
41 {
42  int Nst_pad = CEIL_NWP(Nst);
43 
44  int ngst = NDF * Nst_pad;
45  int up = NDF * Nst_pad * exu;
46  int vp = NDF * Nst_pad * exv;
47  int wp = NDF * Nst_pad * exw;
48 
49 #pragma acc data present(u[up:ngst], v[vp:ngst], w[wp:ngst]) \
50  copyin(a, Nst, Nst_pad, exu, exv, exw)
51 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
52  {
53 
54 #pragma acc loop gang worker vector
55  for(int ist = 0; ist < Nst; ++ist){
56  int igw = ist + Nst_pad * exw;
57  int igv = ist + Nst_pad * exv;
58  int igu = ist + Nst_pad * exu;
59 
60 #ifdef SU3_3RD_ROW_RECONST
61  real_t vt[NDF], wt[NVC];
62  for(int ic2 = 0; ic2 < NC-1; ++ic2){
63  for(int ic1 = 0; ic1 < NC; ++ic1){
64  int icr = 2*ic1 + NVC * ic2;
65  int ici = 2*ic1+1 + NVC * ic2;
66  vt[icr] = v[IDX2_G_R(ic1, ic2, igv)];
67  vt[ici] = v[IDX2_G_I(ic1, ic2, igv)];
68  }
69  }
70 
71  vt[12] = EXT_IMG_R(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
72  vt[ 8], vt[ 9], vt[10], vt[11]);
73  vt[13] = EXT_IMG_I(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
74  vt[ 8], vt[ 9], vt[10], vt[11]);
75  vt[14] = EXT_IMG_R(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
76  vt[10], vt[11], vt[ 6], vt[ 7]);
77  vt[15] = EXT_IMG_I(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
78  vt[10], vt[11], vt[ 6], vt[ 7]);
79  vt[16] = EXT_IMG_R(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
80  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
81  vt[17] = EXT_IMG_I(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
82  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
83 #else
84  real_t vt[NDF], wt[NVC];
85  for(int ic2 = 0; ic2 < NC; ++ic2){
86  for(int ic1 = 0; ic1 < NC; ++ic1){
87  int icr = 2*ic1 + NVC * ic2;
88  int ici = 2*ic1+1 + NVC * ic2;
89  vt[icr] = v[IDX2_G_R(ic1, ic2, igv)];
90  vt[ici] = v[IDX2_G_I(ic1, ic2, igv)];
91  }
92  }
93 #endif
94 
95  for(int ic2 = 0; ic2 < NC; ++ic2){
96 
97  for(int ic1 = 0; ic1 < NC; ++ic1){
98  int icr = 2*ic1;
99  int ici = 2*ic1+1;
100  wt[icr] = w[IDX2_G_R(ic2, ic1, igw)];
101  wt[ici] = w[IDX2_G_I(ic2, ic1, igw)];
102  }
103 
104  for(int ic1 = 0; ic1 < NC; ++ic1){
105  int j = NVC * ic1;
106  real_t ur, ui;
107  ur = MULT_GXr(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
108  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
109  ui = MULT_GXi(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
110  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
111  u[IDX2_G_R(ic2, ic1, igu)] += a * ur;
112  u[IDX2_G_I(ic2, ic1, igu)] += a * ui;
113  }
114 
115  }
116 
117  }
118 
119  } // acc parallel
120 
121 }
122 
123 //====================================================================
124 void mult_Gnn(real_t *RESTRICT u, const int exu,
125  real_t *RESTRICT v, const int exv,
126  real_t *RESTRICT w, const int exw, const int Nst)
127 {
128  int Nst_pad = CEIL_NWP(Nst);
129 
130  int ngst = NDF * Nst_pad;
131  int up = NDF * Nst_pad * exu;
132  int vp = NDF * Nst_pad * exv;
133  int wp = NDF * Nst_pad * exw;
134 
135 #pragma acc data present(u[up:ngst], v[vp:ngst], w[wp:ngst]) \
136  copyin(Nst, Nst_pad, exu, exv, exw)
137 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
138  {
139 
140 #pragma acc loop gang worker vector
141  for(int ist = 0; ist < Nst; ++ist){
142  int igw = ist + Nst_pad * exw;
143  int igv = ist + Nst_pad * exv;
144  int igu = ist + Nst_pad * exu;
145 
146 #ifdef SU3_3RD_ROW_RECONST
147  real_t vt[NDF], wt[NVC];
148  for(int ic2 = 0; ic2 < NC-1; ++ic2){
149  for(int ic1 = 0; ic1 < NC; ++ic1){
150  int icr = 2*ic1 + NVC * ic2;
151  int ici = 2*ic1+1 + NVC * ic2;
152  vt[icr] = v[IDX2_G_R(ic1, ic2, igv)];
153  vt[ici] = v[IDX2_G_I(ic1, ic2, igv)];
154  }
155  }
156 
157  vt[12] = EXT_IMG_R(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
158  vt[ 8], vt[ 9], vt[10], vt[11]);
159  vt[13] = EXT_IMG_I(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
160  vt[ 8], vt[ 9], vt[10], vt[11]);
161  vt[14] = EXT_IMG_R(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
162  vt[10], vt[11], vt[ 6], vt[ 7]);
163  vt[15] = EXT_IMG_I(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
164  vt[10], vt[11], vt[ 6], vt[ 7]);
165  vt[16] = EXT_IMG_R(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
166  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
167  vt[17] = EXT_IMG_I(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
168  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
169 #else
170  real_t vt[NDF], wt[NVC];
171  for(int ic2 = 0; ic2 < NC; ++ic2){
172  for(int ic1 = 0; ic1 < NC; ++ic1){
173  int icr = 2*ic1 + NVC * ic2;
174  int ici = 2*ic1+1 + NVC * ic2;
175  vt[icr] = v[IDX2_G_R(ic1, ic2, igv)];
176  vt[ici] = v[IDX2_G_I(ic1, ic2, igv)];
177  }
178  }
179 #endif
180 
181  for(int ic2 = 0; ic2 < NC; ++ic2){
182 
183  for(int ic1 = 0; ic1 < NC; ++ic1){
184  int icr = 2*ic1;
185  int ici = 2*ic1+1;
186  wt[icr] = w[IDX2_G_R(ic2, ic1, igw)];
187  wt[ici] = w[IDX2_G_I(ic2, ic1, igw)];
188  }
189 
190  for(int ic1 = 0; ic1 < NC; ++ic1){
191  int j = NVC * ic1;
192  real_t ur, ui;
193  ur = MULT_GXr(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
194  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
195  ui = MULT_GXi(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
196  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
197  u[IDX2_G_R(ic2, ic1, igu)] = ur;
198  u[IDX2_G_I(ic2, ic1, igu)] = ui;
199  }
200 
201  }
202 
203  }
204 
205  } // acc parallel
206 
207 }
208 
209 //====================================================================
210 void multadd_Gnd(real_t *RESTRICT u, const int exu,
211  real_t *RESTRICT v, const int exv,
212  real_t *RESTRICT w, const int exw,
213  const real_t a, const int Nst)
214 {
215  int Nst_pad = CEIL_NWP(Nst);
216 
217  int ngst = NDF * Nst_pad;
218  int up = NDF * Nst_pad * exu;
219  int vp = NDF * Nst_pad * exv;
220  int wp = NDF * Nst_pad * exw;
221 
222 #pragma acc data present(u[up:ngst], v[vp:ngst], w[wp:ngst]) \
223  copyin(a, Nst, Nst_pad, exu, exv, exw)
224 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
225  {
226 
227 #pragma acc loop gang worker vector
228  for(int ist = 0; ist < Nst; ++ist){
229  int igw = ist + Nst_pad * exw;
230  int igv = ist + Nst_pad * exv;
231  int igu = ist + Nst_pad * exu;
232 
233 #ifdef SU3_3RD_ROW_RECONST
234  real_t vt[NDF], wt[NVC];
235  for(int ic2 = 0; ic2 < NC-1; ++ic2){
236  for(int ic1 = 0; ic1 < NC; ++ic1){
237  int icr = 2*ic1 + NVC * ic2;
238  int ici = 2*ic1+1 + NVC * ic2;
239  vt[icr] = v[IDX2_G_R(ic1, ic2, igv)];
240  vt[ici] = v[IDX2_G_I(ic1, ic2, igv)];
241  }
242  }
243 
244  vt[12] = EXT_IMG_R(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
245  vt[ 8], vt[ 9], vt[10], vt[11]);
246  vt[13] = EXT_IMG_I(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
247  vt[ 8], vt[ 9], vt[10], vt[11]);
248  vt[14] = EXT_IMG_R(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
249  vt[10], vt[11], vt[ 6], vt[ 7]);
250  vt[15] = EXT_IMG_I(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
251  vt[10], vt[11], vt[ 6], vt[ 7]);
252  vt[16] = EXT_IMG_R(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
253  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
254  vt[17] = EXT_IMG_I(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
255  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
256 #else
257  real_t vt[NDF], wt[NVC];
258  for(int ic2 = 0; ic2 < NC; ++ic2){
259  for(int ic1 = 0; ic1 < NC; ++ic1){
260  int icr = 2*ic1 + NVC * ic2;
261  int ici = 2*ic1+1 + NVC * ic2;
262  vt[icr] = v[IDX2_G_R(ic1, ic2, igv)];
263  vt[ici] = v[IDX2_G_I(ic1, ic2, igv)];
264  }
265  }
266 #endif
267 
268  for(int ic2 = 0; ic2 < NC; ++ic2){
269 
270  for(int ic1 = 0; ic1 < NC; ++ic1){
271  int icr = 2*ic1;
272  int ici = 2*ic1+1;
273  wt[icr] = w[IDX2_G_R(ic1, ic2, igw)];
274  wt[ici] = -w[IDX2_G_I(ic1, ic2, igw)];
275  }
276 
277  for(int ic1 = 0; ic1 < NC; ++ic1){
278  int j = NVC * ic1;
279  real_t ur, ui;
280  ur = MULT_GXr(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
281  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
282  ui = MULT_GXi(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
283  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
284  u[IDX2_G_R(ic2, ic1, igu)] += a * ur;
285  u[IDX2_G_I(ic2, ic1, igu)] += a * ui;
286  }
287 
288  }
289 
290  }
291 
292  } // acc parallel
293 
294 }
295 
296 //====================================================================
297 void mult_Gnd(real_t *RESTRICT u, const int exu,
298  real_t *RESTRICT v, const int exv,
299  real_t *RESTRICT w, const int exw, const int Nst)
300 {
301  int Nst_pad = CEIL_NWP(Nst);
302 
303  int ngst = NDF * Nst_pad;
304  int up = NDF * Nst_pad * exu;
305  int vp = NDF * Nst_pad * exv;
306  int wp = NDF * Nst_pad * exw;
307 
308 #pragma acc data present(u[up:ngst], v[vp:ngst], w[wp:ngst]) \
309  copyin(Nst, Nst_pad, exu, exv, exw)
310 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
311  {
312 
313 #pragma acc loop gang worker vector
314  for(int ist = 0; ist < Nst; ++ist){
315  int igw = ist + Nst_pad * exw;
316  int igv = ist + Nst_pad * exv;
317  int igu = ist + Nst_pad * exu;
318 
319 #ifdef SU3_3RD_ROW_RECONST
320  real_t vt[NDF], wt[NVC];
321  for(int ic2 = 0; ic2 < NC-1; ++ic2){
322  for(int ic1 = 0; ic1 < NC; ++ic1){
323  int icr = 2*ic1 + NVC * ic2;
324  int ici = 2*ic1+1 + NVC * ic2;
325  vt[icr] = v[IDX2_G_R(ic1, ic2, igv)];
326  vt[ici] = v[IDX2_G_I(ic1, ic2, igv)];
327  }
328  }
329 
330  vt[12] = EXT_IMG_R(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
331  vt[ 8], vt[ 9], vt[10], vt[11]);
332  vt[13] = EXT_IMG_I(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
333  vt[ 8], vt[ 9], vt[10], vt[11]);
334  vt[14] = EXT_IMG_R(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
335  vt[10], vt[11], vt[ 6], vt[ 7]);
336  vt[15] = EXT_IMG_I(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
337  vt[10], vt[11], vt[ 6], vt[ 7]);
338  vt[16] = EXT_IMG_R(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
339  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
340  vt[17] = EXT_IMG_I(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
341  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
342 #else
343  real_t vt[NDF], wt[NVC];
344  for(int ic2 = 0; ic2 < NC; ++ic2){
345  for(int ic1 = 0; ic1 < NC; ++ic1){
346  int icr = 2*ic1 + NVC * ic2;
347  int ici = 2*ic1+1 + NVC * ic2;
348  vt[icr] = v[IDX2_G_R(ic1, ic2, igv)];
349  vt[ici] = v[IDX2_G_I(ic1, ic2, igv)];
350  }
351  }
352 #endif
353 
354  for(int ic2 = 0; ic2 < NC; ++ic2){
355 
356  for(int ic1 = 0; ic1 < NC; ++ic1){
357  int icr = 2*ic1;
358  int ici = 2*ic1+1;
359  wt[icr] = w[IDX2_G_R(ic1, ic2, igw)];
360  wt[ici] = -w[IDX2_G_I(ic1, ic2, igw)];
361  }
362 
363  for(int ic1 = 0; ic1 < NC; ++ic1){
364  int j = NVC * ic1;
365  real_t ur, ui;
366  ur = MULT_GXr(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
367  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
368  ui = MULT_GXi(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
369  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
370  u[IDX2_G_R(ic2, ic1, igu)] = ur;
371  u[IDX2_G_I(ic2, ic1, igu)] = ui;
372  }
373 
374  }
375 
376  }
377 
378  } // acc parallel
379 
380 }
381 
382 //====================================================================
383 void multadd_Gdn(real_t *RESTRICT u, const int exu,
384  real_t *RESTRICT v, const int exv,
385  real_t *RESTRICT w, const int exw,
386  const real_t a, const int Nst)
387 {
388  int Nst_pad = CEIL_NWP(Nst);
389 
390  int ngst = NDF * Nst_pad;
391  int up = NDF * Nst_pad * exu;
392  int vp = NDF * Nst_pad * exv;
393  int wp = NDF * Nst_pad * exw;
394 
395 #pragma acc data present(u[up:ngst], v[vp:ngst], w[wp:ngst]) \
396  copyin(a, Nst, Nst_pad, exu, exv, exw)
397 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
398  {
399 
400 #pragma acc loop gang worker vector
401  for(int ist = 0; ist < Nst; ++ist){
402  int igw = ist + Nst_pad * exw;
403  int igv = ist + Nst_pad * exv;
404  int igu = ist + Nst_pad * exu;
405 
406 #ifdef SU3_3RD_ROW_RECONST
407  real_t vt[NDF], wt[NVC];
408  for(int ic2 = 0; ic2 < NC-1; ++ic2){
409  for(int ic1 = 0; ic1 < NC; ++ic1){
410  int icr = 2*ic1 + NVC * ic2;
411  int ici = 2*ic1+1 + NVC * ic2;
412  vt[icr] = v[IDX2_G_R(ic2, ic1, igv)];
413  vt[ici] = -v[IDX2_G_I(ic2, ic1, igv)];
414  }
415  }
416 
417  vt[12] = EXT_IMG_R(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
418  vt[ 8], vt[ 9], vt[10], vt[11]);
419  vt[13] = EXT_IMG_I(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
420  vt[ 8], vt[ 9], vt[10], vt[11]);
421  vt[14] = EXT_IMG_R(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
422  vt[10], vt[11], vt[ 6], vt[ 7]);
423  vt[15] = EXT_IMG_I(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
424  vt[10], vt[11], vt[ 6], vt[ 7]);
425  vt[16] = EXT_IMG_R(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
426  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
427  vt[17] = EXT_IMG_I(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
428  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
429 #else
430  real_t vt[NDF], wt[NVC];
431  for(int ic2 = 0; ic2 < NC; ++ic2){
432  for(int ic1 = 0; ic1 < NC; ++ic1){
433  int icr = 2*ic1 + NVC * ic2;
434  int ici = 2*ic1+1 + NVC * ic2;
435  vt[icr] = v[IDX2_G_R(ic2, ic1, igv)];
436  vt[ici] = -v[IDX2_G_I(ic2, ic1, igv)];
437  }
438  }
439 #endif
440 
441  for(int ic2 = 0; ic2 < NC; ++ic2){
442 
443  for(int ic1 = 0; ic1 < NC; ++ic1){
444  int icr = 2*ic1;
445  int ici = 2*ic1+1;
446  wt[icr] = w[IDX2_G_R(ic2, ic1, igw)];
447  wt[ici] = w[IDX2_G_I(ic2, ic1, igw)];
448  }
449 
450  for(int ic1 = 0; ic1 < NC; ++ic1){
451  int j = NVC * ic1;
452  real_t ur, ui;
453  ur = MULT_GXr(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
454  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
455  ui = MULT_GXi(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
456  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
457  u[IDX2_G_R(ic2, ic1, igu)] += a * ur;
458  u[IDX2_G_I(ic2, ic1, igu)] += a * ui;
459  }
460 
461  }
462 
463  }
464 
465  } // acc parallel
466 
467 }
468 
469 //====================================================================
470 void mult_Gdn(real_t *RESTRICT u, const int exu,
471  real_t *RESTRICT v, const int exv,
472  real_t *RESTRICT w, const int exw, const int Nst)
473 {
474  int Nst_pad = CEIL_NWP(Nst);
475 
476  int ngst = NDF * Nst_pad;
477  int up = NDF * Nst_pad * exu;
478  int vp = NDF * Nst_pad * exv;
479  int wp = NDF * Nst_pad * exw;
480 
481 #pragma acc data present(u[up:ngst], v[vp:ngst], w[wp:ngst]) \
482  copyin(Nst, Nst_pad, exu, exv, exw)
483 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
484  {
485 
486 #pragma acc loop gang worker vector
487  for(int ist = 0; ist < Nst; ++ist){
488  int igw = ist + Nst_pad * exw;
489  int igv = ist + Nst_pad * exv;
490  int igu = ist + Nst_pad * exu;
491 
492 #ifdef SU3_3RD_ROW_RECONST
493  real_t vt[NDF], wt[NVC];
494  for(int ic2 = 0; ic2 < NC-1; ++ic2){
495  for(int ic1 = 0; ic1 < NC; ++ic1){
496  int icr = 2*ic1 + NVC * ic2;
497  int ici = 2*ic1+1 + NVC * ic2;
498  vt[icr] = v[IDX2_G_R(ic2, ic1, igv)];
499  vt[ici] = -v[IDX2_G_I(ic2, ic1, igv)];
500  }
501  }
502 
503  vt[12] = EXT_IMG_R(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
504  vt[ 8], vt[ 9], vt[10], vt[11]);
505  vt[13] = EXT_IMG_I(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
506  vt[ 8], vt[ 9], vt[10], vt[11]);
507  vt[14] = EXT_IMG_R(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
508  vt[10], vt[11], vt[ 6], vt[ 7]);
509  vt[15] = EXT_IMG_I(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
510  vt[10], vt[11], vt[ 6], vt[ 7]);
511  vt[16] = EXT_IMG_R(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
512  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
513  vt[17] = EXT_IMG_I(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
514  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
515 #else
516  real_t vt[NDF], wt[NVC];
517  for(int ic2 = 0; ic2 < NC; ++ic2){
518  for(int ic1 = 0; ic1 < NC; ++ic1){
519  int icr = 2*ic1 + NVC * ic2;
520  int ici = 2*ic1+1 + NVC * ic2;
521  vt[icr] = v[IDX2_G_R(ic2, ic1, igv)];
522  vt[ici] = -v[IDX2_G_I(ic2, ic1, igv)];
523  }
524  }
525 #endif
526 
527  for(int ic2 = 0; ic2 < NC; ++ic2){
528 
529  for(int ic1 = 0; ic1 < NC; ++ic1){
530  int icr = 2*ic1;
531  int ici = 2*ic1+1;
532  wt[icr] = w[IDX2_G_R(ic2, ic1, igw)];
533  wt[ici] = w[IDX2_G_I(ic2, ic1, igw)];
534  }
535 
536  for(int ic1 = 0; ic1 < NC; ++ic1){
537  int j = NVC * ic1;
538  real_t ur, ui;
539  ur = MULT_GXr(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
540  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
541  ui = MULT_GXi(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
542  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
543  u[IDX2_G_R(ic2, ic1, igu)] = ur;
544  u[IDX2_G_I(ic2, ic1, igu)] = ui;
545  }
546 
547  }
548 
549  }
550 
551  } // acc parallel
552 
553 }
554 
555 //====================================================================
556 void mult_Gdd(real_t *RESTRICT u, const int exu,
557  real_t *RESTRICT v, const int exv,
558  real_t *RESTRICT w, const int exw, const int Nst)
559 {
560  int Nst_pad = CEIL_NWP(Nst);
561 
562  int ngst = NDF * Nst_pad;
563  int up = NDF * Nst_pad * exu;
564  int vp = NDF * Nst_pad * exv;
565  int wp = NDF * Nst_pad * exw;
566 
567 #pragma acc data present(u[up:ngst], v[vp:ngst], w[wp:ngst]) \
568  copyin(Nst, Nst_pad, exu, exv, exw)
569 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
570  {
571 
572 #pragma acc loop gang worker vector
573  for(int ist = 0; ist < Nst; ++ist){
574  int igw = ist + Nst_pad * exw;
575  int igv = ist + Nst_pad * exv;
576  int igu = ist + Nst_pad * exu;
577 
578 #ifdef SU3_3RD_ROW_RECONST
579  real_t vt[NDF], wt[NVC];
580  for(int ic2 = 0; ic2 < NC-1; ++ic2){
581  for(int ic1 = 0; ic1 < NC; ++ic1){
582  int icr = 2*ic1 + NVC * ic2;
583  int ici = 2*ic1+1 + NVC * ic2;
584  vt[icr] = v[IDX2_G_R(ic2, ic1, igv)];
585  vt[ici] = -v[IDX2_G_I(ic2, ic1, igv)];
586  }
587  }
588 
589  vt[12] = EXT_IMG_R(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
590  vt[ 8], vt[ 9], vt[10], vt[11]);
591  vt[13] = EXT_IMG_I(vt[ 2], vt[ 3], vt[ 4], vt[ 5],
592  vt[ 8], vt[ 9], vt[10], vt[11]);
593  vt[14] = EXT_IMG_R(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
594  vt[10], vt[11], vt[ 6], vt[ 7]);
595  vt[15] = EXT_IMG_I(vt[ 4], vt[ 5], vt[ 0], vt[ 1],
596  vt[10], vt[11], vt[ 6], vt[ 7]);
597  vt[16] = EXT_IMG_R(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
598  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
599  vt[17] = EXT_IMG_I(vt[ 0], vt[ 1], vt[ 2], vt[ 3],
600  vt[ 6], vt[ 7], vt[ 8], vt[ 9]);
601 #else
602  real_t vt[NDF], wt[NVC];
603  for(int ic2 = 0; ic2 < NC; ++ic2){
604  for(int ic1 = 0; ic1 < NC; ++ic1){
605  int icr = 2*ic1 + NVC * ic2;
606  int ici = 2*ic1+1 + NVC * ic2;
607  vt[icr] = v[IDX2_G_R(ic2, ic1, igv)];
608  vt[ici] = -v[IDX2_G_I(ic2, ic1, igv)];
609  }
610  }
611 #endif
612 
613  for(int ic2 = 0; ic2 < NC; ++ic2){
614 
615  for(int ic1 = 0; ic1 < NC; ++ic1){
616  int icr = 2*ic1;
617  int ici = 2*ic1+1;
618  wt[icr] = w[IDX2_G_R(ic1, ic2, igw)];
619  wt[ici] = -w[IDX2_G_I(ic1, ic2, igw)];
620  }
621 
622  for(int ic1 = 0; ic1 < NC; ++ic1){
623  int j = NVC * ic1;
624  real_t ur, ui;
625  ur = MULT_GXr(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
626  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
627  ui = MULT_GXi(vt[0+j], vt[1+j], vt[2+j], vt[3+j], vt[4+j], vt[5+j],
628  wt[0], wt[1], wt[2], wt[3], wt[4], wt[5]);
629  u[IDX2_G_R(ic2, ic1, igu)] = ur;
630  u[IDX2_G_I(ic2, ic1, igu)] = ui;
631  }
632 
633  }
634 
635  }
636 
637  } // acc parallel
638 
639 }
640 
641 //====================================================================
642 // anti-hermitian
643 void ah_G(real_t *u, const int ex, const int Nst)
644 {
645  int Nst_pad = CEIL_NWP(Nst);
646 
647  int ngst = NDF * Nst_pad;
648  int up = NDF * Nst_pad * ex;
649 
650 #pragma acc data present(u[up:ngst]) copyin(ex, Nst, Nst_pad)
651 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
652  {
653 
654 #pragma acc loop gang worker vector
655  for(int ist = 0; ist < Nst; ++ist){
656  int igu = ist + Nst_pad * ex;
657 
658  real_t ut[NDF];
659  for(int ic2 = 0; ic2 < NC; ++ic2){
660  for(int ic1 = 0; ic1 < NC; ++ic1){
661  int icr = 2*ic1 + NVC * ic2;
662  int ici = 2*ic1+1 + NVC * ic2;
663  ut[icr] = u[IDX2_G_R(ic1, ic2, igu)];
664  ut[ici] = u[IDX2_G_I(ic1, ic2, igu)];
665  }
666  }
667 
668  real_t vt[NDF];
669  vt[ 0] = 0.0;
670  vt[ 1] = ut[ 1];
671  vt[ 2] = 0.5 *(ut[ 2] - ut[ 6]);
672  vt[ 3] = 0.5 *(ut[ 3] + ut[ 7]);
673  vt[ 4] = 0.5 *(ut[ 4] - ut[12]);
674  vt[ 5] = 0.5 *(ut[ 5] + ut[13]);
675  vt[ 6] = -vt[ 2];
676  vt[ 7] = vt[ 3];
677  vt[ 8] = 0.0;
678  vt[ 9] = ut[ 9];
679  vt[10] = 0.5 *(ut[10] - ut[14]);
680  vt[11] = 0.5 *(ut[11] + ut[15]);
681  vt[12] = -vt[ 4];
682  vt[13] = vt[ 5];
683  vt[14] = -vt[10];
684  vt[15] = vt[11];
685  vt[16] = 0.0;
686  vt[17] = ut[17];
687 
688  for(int ic2 = 0; ic2 < NC; ++ic2){
689  for(int ic1 = 0; ic1 < NC; ++ic1){
690  int icr = 2*ic1 + NVC * ic2;
691  int ici = 2*ic1+1 + NVC * ic2;
692  u[IDX2_G_R(ic1, ic2, igu)] = vt[icr];
693  u[IDX2_G_I(ic1, ic2, igu)] = vt[ici];
694  }
695  }
696 
697  }
698 
699  } // acc parallel
700 
701 }
702 
703 //====================================================================
704 // anti-hermitian traceless
705 void at_G(real_t *u, const int ex, const int Nst)
706 {
707  int Nst_pad = CEIL_NWP(Nst);
708 
709  int ngst = NDF * Nst_pad;
710  int up = NDF * Nst_pad * ex;
711 
712 #pragma acc data present(u[up:ngst]) copyin(ex, Nst, Nst_pad)
713 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
714  {
715 
716 #pragma acc loop gang worker vector
717  for(int ist = 0; ist < Nst; ++ist){
718  int igu = ist + Nst_pad * ex;
719 
720  real_t ut[NDF];
721  for(int ic2 = 0; ic2 < NC; ++ic2){
722  for(int ic1 = 0; ic1 < NC; ++ic1){
723  int icr = 2*ic1 + NVC * ic2;
724  int ici = 2*ic1+1 + NVC * ic2;
725  ut[icr] = u[IDX2_G_R(ic1, ic2, igu)];
726  ut[ici] = u[IDX2_G_I(ic1, ic2, igu)];
727  }
728  }
729 
730  real_t vt[NDF];
731  real_t tr = (ut[ 1] + ut[9] + ut[17])/3.0;
732  vt[ 0] = 0.0;
733  vt[ 1] = ut[ 1] - tr;
734  vt[ 2] = 0.5 *(ut[ 2] - ut[ 6]);
735  vt[ 3] = 0.5 *(ut[ 3] + ut[ 7]);
736  vt[ 4] = 0.5 *(ut[ 4] - ut[12]);
737  vt[ 5] = 0.5 *(ut[ 5] + ut[13]);
738  vt[ 6] = -vt[ 2];
739  vt[ 7] = vt[ 3];
740  vt[ 8] = 0.0;
741  vt[ 9] = ut[ 9] - tr;
742  vt[10] = 0.5 *(ut[10] - ut[14]);
743  vt[11] = 0.5 *(ut[11] + ut[15]);
744  vt[12] = -vt[ 4];
745  vt[13] = vt[ 5];
746  vt[14] = -vt[10];
747  vt[15] = vt[11];
748  vt[16] = 0.0;
749  vt[17] = ut[17] - tr;
750 
751  for(int ic2 = 0; ic2 < NC; ++ic2){
752  for(int ic1 = 0; ic1 < NC; ++ic1){
753  int icr = 2*ic1 + NVC * ic2;
754  int ici = 2*ic1+1 + NVC * ic2;
755  u[IDX2_G_R(ic1, ic2, igu)] = vt[icr];
756  u[IDX2_G_I(ic1, ic2, igu)] = vt[ici];
757  }
758  }
759 
760  }
761 
762  } // acc parallel
763 
764 }
765 
766 //====================================================================
767 void add_unit(real_t *u, const int ex, const real_t a, const int Nst)
768  // u = u + a * I (I: unit matrix)
769 {
770  int Nst_pad = CEIL_NWP(Nst);
771 
772  int ngst = NDF * Nst_pad;
773  int up = NDF * Nst_pad * ex;
774 
775 #pragma acc data present(u[up:ngst]) copyin(a, ex, Nst, Nst_pad)
776 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
777  {
778 
779 #pragma acc loop gang worker vector
780  for(int ist = 0; ist < Nst; ++ist){
781  int igu = ist + Nst_pad * ex;
782 
783  for(int ic = 0; ic < NC; ++ic){
784  real_t ut = u[IDX2_G_R(ic, ic, igu)];
785  ut += a;
786  u[IDX2_G_R(ic, ic, igu)] = ut;
787  }
788 
789  }
790 
791  } // acc parallel
792 
793 }
794 
795 //====================================================================
796 void inverse_dag(real_t *RESTRICT u, const int ex1,
797  const real_t *RESTRICT v, const int ex2,
798  const int Nst)
799  // calculate u = ((v)^inv)^dag
800 {
801  int Nst_pad = CEIL_NWP(Nst);
802 
803  int ngst = NDF * Nst_pad;
804  int up = NDF * Nst_pad * ex1;
805  int vp = NDF * Nst_pad * ex2;
806 
807 #pragma acc data present(v[vp:ngst], u[up:ngst]) \
808  copyin(ex1, ex2, Nst, Nst_pad)
809 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
810  {
811 
812 #pragma acc loop gang worker vector
813  for(int ist = 0; ist < Nst; ++ist){
814 
815  int ist1 = ist + Nst_pad * ex1;
816  int ist2 = ist + Nst_pad * ex2;
817 
818  real_t vt[18], wt[18], ut[18];
819 
820  vt[ 0] = v[IDX2_G_R(0, 0, ist2)];
821  vt[ 1] = v[IDX2_G_I(0, 0, ist2)];
822  vt[ 2] = v[IDX2_G_R(1, 0, ist2)];
823  vt[ 3] = v[IDX2_G_I(1, 0, ist2)];
824  vt[ 4] = v[IDX2_G_R(2, 0, ist2)];
825  vt[ 5] = v[IDX2_G_I(2, 0, ist2)];
826 
827  vt[ 6] = v[IDX2_G_R(0, 1, ist2)];
828  vt[ 7] = v[IDX2_G_I(0, 1, ist2)];
829  vt[ 8] = v[IDX2_G_R(1, 1, ist2)];
830  vt[ 9] = v[IDX2_G_I(1, 1, ist2)];
831  vt[10] = v[IDX2_G_R(2, 1, ist2)];
832  vt[11] = v[IDX2_G_I(2, 1, ist2)];
833 
834  vt[12] = v[IDX2_G_R(0, 2, ist2)];
835  vt[13] = v[IDX2_G_I(0, 2, ist2)];
836  vt[14] = v[IDX2_G_R(1, 2, ist2)];
837  vt[15] = v[IDX2_G_I(1, 2, ist2)];
838  vt[16] = v[IDX2_G_R(2, 2, ist2)];
839  vt[17] = v[IDX2_G_I(2, 2, ist2)];
840 
841  wt[ 0] = MULT_R(vt[ 8], vt[ 9], vt[16], vt[17])
842  - MULT_R(vt[10], vt[11], vt[14], vt[15]);
843  wt[ 1] = MULT_I(vt[ 8], vt[ 9], vt[16], vt[17])
844  - MULT_I(vt[10], vt[11], vt[14], vt[15]);
845 
846  wt[ 2] = MULT_R(vt[10], vt[11], vt[12], vt[13])
847  - MULT_R(vt[ 6], vt[ 7], vt[16], vt[17]);
848  wt[ 3] = MULT_I(vt[10], vt[11], vt[12], vt[13])
849  - MULT_I(vt[ 6], vt[ 7], vt[16], vt[17]);
850 
851  wt[ 4] = MULT_R(vt[ 6], vt[ 7], vt[14], vt[15])
852  - MULT_R(vt[ 8], vt[ 9], vt[12], vt[13]);
853  wt[ 5] = MULT_I(vt[ 6], vt[ 7], vt[14], vt[15])
854  - MULT_I(vt[ 8], vt[ 9], vt[12], vt[13]);
855 
856  wt[ 6] = MULT_R(vt[14], vt[15], vt[ 4], vt[ 5])
857  - MULT_R(vt[16], vt[17], vt[ 2], vt[ 3]);
858  wt[ 7] = MULT_I(vt[14], vt[15], vt[ 4], vt[ 5])
859  - MULT_I(vt[16], vt[17], vt[ 2], vt[ 3]);
860 
861  wt[ 8] = MULT_R(vt[16], vt[17], vt[ 0], vt[ 1])
862  - MULT_R(vt[12], vt[13], vt[ 4], vt[ 5]);
863  wt[ 9] = MULT_I(vt[16], vt[17], vt[ 0], vt[ 1])
864  - MULT_I(vt[12], vt[13], vt[ 4], vt[ 5]);
865 
866  wt[10] = MULT_R(vt[12], vt[13], vt[ 2], vt[ 3])
867  - MULT_R(vt[14], vt[15], vt[ 0], vt[ 1]);
868  wt[11] = MULT_I(vt[12], vt[13], vt[ 2], vt[ 3])
869  - MULT_I(vt[14], vt[15], vt[ 0], vt[ 1]);
870 
871  wt[12] = MULT_R(vt[ 2], vt[ 3], vt[10], vt[11])
872  - MULT_R(vt[ 4], vt[ 5], vt[ 8], vt[ 9]);
873  wt[13] = MULT_I(vt[ 2], vt[ 3], vt[10], vt[11])
874  - MULT_I(vt[ 4], vt[ 5], vt[ 8], vt[ 9]);
875 
876  wt[14] = MULT_R(vt[ 4], vt[ 5], vt[ 6], vt[ 7])
877  - MULT_R(vt[ 0], vt[ 1], vt[10], vt[11]);
878  wt[15] = MULT_I(vt[ 4], vt[ 5], vt[ 6], vt[ 7])
879  - MULT_I(vt[ 0], vt[ 1], vt[10], vt[11]);
880 
881  wt[16] = MULT_R(vt[ 0], vt[ 1], vt[ 8], vt[ 9])
882  - MULT_R(vt[ 2], vt[ 3], vt[ 6], vt[ 7]);
883  wt[17] = MULT_I(vt[ 0], vt[ 1], vt[ 8], vt[ 9])
884  - MULT_I(vt[ 2], vt[ 3], vt[ 6], vt[ 7]);
885 
886  real_t det_r = MULT_R(vt[0], vt[1], wt[0], wt[1])
887  + MULT_R(vt[2], vt[3], wt[2], wt[3])
888  + MULT_R(vt[4], vt[5], wt[4], wt[5]);
889  real_t det_i = MULT_I(vt[0], vt[1], wt[0], wt[1])
890  + MULT_I(vt[2], vt[3], wt[2], wt[3])
891  + MULT_I(vt[4], vt[5], wt[4], wt[5]);
892 
893  real_t det2 = det_r * det_r + det_i * det_i;
894  real_t detinv_r = det_r/det2;
895  real_t detinv_i = - det_i/det2;
896 
897  ut[ 0] = MULT_R(wt[ 0], wt[ 1], detinv_r, detinv_i);
898  ut[ 1] = MULT_I(wt[ 0], wt[ 1], detinv_r, detinv_i);
899 
900  ut[ 2] = MULT_R(wt[ 6], wt[ 7], detinv_r, detinv_i);
901  ut[ 3] = MULT_I(wt[ 6], wt[ 7], detinv_r, detinv_i);
902 
903  ut[ 4] = MULT_R(wt[12], wt[13], detinv_r, detinv_i);
904  ut[ 5] = MULT_I(wt[12], wt[13], detinv_r, detinv_i);
905 
906  ut[ 6] = MULT_R(wt[ 2], wt[ 3], detinv_r, detinv_i);
907  ut[ 7] = MULT_I(wt[ 2], wt[ 3], detinv_r, detinv_i);
908 
909  ut[ 8] = MULT_R(wt[ 8], wt[ 9], detinv_r, detinv_i);
910  ut[ 9] = MULT_I(wt[ 8], wt[ 9], detinv_r, detinv_i);
911 
912  ut[10] = MULT_R(wt[14], wt[15], detinv_r, detinv_i);
913  ut[11] = MULT_I(wt[14], wt[15], detinv_r, detinv_i);
914 
915  ut[12] = MULT_R(wt[ 4], wt[ 5], detinv_r, detinv_i);
916  ut[13] = MULT_I(wt[ 4], wt[ 5], detinv_r, detinv_i);
917 
918  ut[14] = MULT_R(wt[10], wt[11], detinv_r, detinv_i);
919  ut[15] = MULT_I(wt[10], wt[11], detinv_r, detinv_i);
920 
921  ut[16] = MULT_R(wt[16], wt[17], detinv_r, detinv_i);
922  ut[17] = MULT_I(wt[16], wt[17], detinv_r, detinv_i);
923 
924  u[IDX2_G_R(0, 0, ist1)] = ut[ 0];
925  u[IDX2_G_I(0, 0, ist1)] = -ut[ 1];
926  u[IDX2_G_R(1, 0, ist1)] = ut[ 6];
927  u[IDX2_G_I(1, 0, ist1)] = -ut[ 7];
928  u[IDX2_G_R(2, 0, ist1)] = ut[12];
929  u[IDX2_G_I(2, 0, ist1)] = -ut[13];
930 
931  u[IDX2_G_R(0, 1, ist1)] = ut[ 2];
932  u[IDX2_G_I(0, 1, ist1)] = -ut[ 3];
933  u[IDX2_G_R(1, 1, ist1)] = ut[ 8];
934  u[IDX2_G_I(1, 1, ist1)] = -ut[ 9];
935  u[IDX2_G_R(2, 1, ist1)] = ut[14];
936  u[IDX2_G_I(2, 1, ist1)] = -ut[15];
937 
938  u[IDX2_G_R(0, 2, ist1)] = ut[ 4];
939  u[IDX2_G_I(0, 2, ist1)] = -ut[ 5];
940  u[IDX2_G_R(1, 2, ist1)] = ut[10];
941  u[IDX2_G_I(1, 2, ist1)] = -ut[11];
942  u[IDX2_G_R(2, 2, ist1)] = ut[16];
943  u[IDX2_G_I(2, 2, ist1)] = -ut[17];
944 
945  }
946 
947  } // acc parallel
948 
949 }
950 
951 //============================================================END=====
952 #endif
ah_G
void ah_G(real_t *u, const int ex, const int Nst)
Definition: afield_Gauge_openacc-inc.h:643
multadd_Gnd
void multadd_Gnd(real_t *RESTRICT u, const int exu, real_t *RESTRICT v, const int exv, real_t *RESTRICT w, const int exw, const real_t a, const int Nst)
Definition: afield_Gauge_openacc-inc.h:210
add_unit
void add_unit(real_t *u, const int ex, const real_t a, const int Nst)
Definition: afield_Gauge_openacc-inc.h:767
multadd_Gdn
void multadd_Gdn(real_t *RESTRICT u, const int exu, real_t *RESTRICT v, const int exv, real_t *RESTRICT w, const int exw, const real_t a, const int Nst)
Definition: afield_Gauge_openacc-inc.h:383
mult_r
real_t mult_r(const real_t v1r, const real_t v1i, const real_t v2r, const real_t v2i)
Definition: afield_Gauge_openacc-inc.h:26
inverse_dag
void inverse_dag(real_t *RESTRICT u, const int ex1, const real_t *RESTRICT v, const int ex2, const int Nst)
Definition: afield_Gauge_openacc-inc.h:796
EXT_IMG_R
#define EXT_IMG_R(v1r, v1i, v2r, v2i, w1r, w1i, w2r, w2i)
Definition: afield_Gauge_openacc-inc.h:19
at_G
void at_G(real_t *u, const int ex, const int Nst)
Definition: afield_Gauge_openacc-inc.h:705
wt
real_t wt[NVCD]
Definition: mult_Clover_csw_chiral_openacc-inc.h:9
mult_Gnd
void mult_Gnd(real_t *RESTRICT u, const int exu, real_t *RESTRICT v, const int exv, real_t *RESTRICT w, const int exw, const int Nst)
Definition: afield_Gauge_openacc-inc.h:297
NDF
#define NDF
Definition: field_F_imp_SU2-inc.h:17
MULT_I
#define MULT_I(v1r, v1i, v2r, v2i)
Definition: afield_Gauge_openacc-inc.h:23
MULT_R
#define MULT_R(v1r, v1i, v2r, v2i)
Definition: afield_Gauge_openacc-inc.h:22
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_GXi
#define MULT_GXi(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: afield_Gauge_openacc-inc.h:17
mult_Gdn
void mult_Gdn(real_t *RESTRICT u, const int exu, real_t *RESTRICT v, const int exv, real_t *RESTRICT w, const int exw, const int Nst)
Definition: afield_Gauge_openacc-inc.h:470
NC
#define NC
Definition: field_F_imp_SU2-inc.h:15
multadd_Gnn
void multadd_Gnn(real_t *RESTRICT u, const int exu, real_t *RESTRICT v, const int exv, real_t *RESTRICT w, const int exw, const real_t a, const int Nst)
Definition: afield_Gauge_openacc-inc.h:37
mult_i
real_t mult_i(const real_t v1r, const real_t v1i, const real_t v2r, const real_t v2i)
Definition: afield_Gauge_openacc-inc.h:31
mult_Gdd
void mult_Gdd(real_t *RESTRICT u, const int exu, real_t *RESTRICT v, const int exv, real_t *RESTRICT w, const int exw, const int Nst)
Definition: afield_Gauge_openacc-inc.h:556
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
define_index.h
NVC
#define NVC
Definition: fopr_Wilson_impl_SU2-inc.h:15
MULT_GXr
#define MULT_GXr(u0, u1, u2, u3, u4, u5, v0, v1, v2, v3, v4, v5)
Definition: afield_Gauge_openacc-inc.h:16
mult_Gnn
void mult_Gnn(real_t *RESTRICT u, const int exu, real_t *RESTRICT v, const int exv, real_t *RESTRICT w, const int exw, const int Nst)
Definition: afield_Gauge_openacc-inc.h:124
IDX2_G_R
#define IDX2_G_R(ic1, ic2, ist)
Definition: define_index.h:51
define_params.h
EXT_IMG_I
#define EXT_IMG_I(v1r, v1i, v2r, v2i, w1r, w1i, w2r, w2i)
Definition: afield_Gauge_openacc-inc.h:20