QUDA  1.0.0
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
clover_reference.cpp
Go to the documentation of this file.
1 #include <stdio.h>
2 #include <stdlib.h>
3 #include <math.h>
4 #include <complex>
5 
6 #include <util_quda.h>
7 #include <test_util.h>
9 #include <blas_reference.h>
10 
11 
19 template <typename sFloat, typename cFloat>
20 void cloverReference(sFloat *out, cFloat *clover, sFloat *in, int parity) {
21  int nSpin = 4;
22  int nColor = 3;
23  int N = nColor * nSpin / 2;
24  int chiralBlock = N + 2*(N-1)*N/2;
25 
26  for (int i=0; i<Vh; i++) {
27  std::complex<sFloat> *In = reinterpret_cast<std::complex<sFloat>*>(&in[i*nSpin*nColor*2]);
28  std::complex<sFloat> *Out = reinterpret_cast<std::complex<sFloat>*>(&out[i*nSpin*nColor*2]);
29 
30  for (int chi=0; chi<nSpin/2; chi++) {
31  cFloat *D = &clover[((parity*Vh + i)*2 + chi)*chiralBlock];
32  std::complex<cFloat> *L = reinterpret_cast<std::complex<cFloat>*>(&D[N]);
33 
34  for (int s_col=0; s_col<nSpin/2; s_col++) { // 2 spins per chiral block
35  for (int c_col=0; c_col<nColor; c_col++) {
36  const int col = s_col * nColor + c_col;
37  const int Col = chi*N + col;
38  Out[Col] = 0.0;
39 
40  for (int s_row=0; s_row<nSpin/2; s_row++) { // 2 spins per chiral block
41  for (int c_row=0; c_row<nColor; c_row++) {
42  const int row = s_row * nColor + c_row;
43  const int Row = chi*N + row;
44 
45  if (row == col) {
46  Out[Col] += D[row] * In[Row];
47  } else if (col < row) {
48  int k = N*(N-1)/2 - (N-col)*(N-col-1)/2 + row - col - 1;
49  Out[Col] += conj(L[k]) * In[Row];
50  } else if (row < col) {
51  int k = N*(N-1)/2 - (N-row)*(N-row-1)/2 + col - row - 1;
52  Out[Col] += L[k] * In[Row];
53  }
54  }
55  }
56 
57  }
58  }
59 
60  }
61 
62  }
63 
64 }
65 
66 void apply_clover(void *out, void *clover, void *in, int parity, QudaPrecision precision) {
67 
68  switch (precision) {
70  cloverReference(static_cast<double*>(out), static_cast<double*>(clover), static_cast<double*>(in), parity);
71  break;
73  cloverReference(static_cast<float*>(out), static_cast<float*>(clover), static_cast<float*>(in), parity);
74  break;
75  default:
76  errorQuda("Unsupported precision %d", precision);
77  }
78 
79 }
80 
81 void clover_dslash(void *out, void **gauge, void *clover, void *in, int parity,
82  int dagger, QudaPrecision precision, QudaGaugeParam &param) {
83  void *tmp = malloc(Vh*spinorSiteSize*precision);
84 
85  wil_dslash(tmp, gauge, in, parity, dagger, precision, param);
86  apply_clover(out, clover, tmp, parity, precision);
87 
88  free(tmp);
89 }
90 
91 // Apply the even-odd preconditioned Wilson-clover operator
92 void clover_matpc(void *out, void **gauge, void *clover, void *clover_inv, void *in, double kappa,
94 
95  double kappa2 = -kappa*kappa;
96  void *tmp = malloc(Vh*spinorSiteSize*precision);
97 
98  switch(matpc_type) {
100  if (!dagger) {
101  wil_dslash(tmp, gauge, in, 1, dagger, precision, gauge_param);
102  apply_clover(out, clover_inv, tmp, 1, precision);
103  wil_dslash(tmp, gauge, out, 0, dagger, precision, gauge_param);
104  apply_clover(out, clover_inv, tmp, 0, precision);
105  } else {
106  apply_clover(tmp, clover_inv, in, 0, precision);
107  wil_dslash(out, gauge, tmp, 1, dagger, precision, gauge_param);
108  apply_clover(tmp, clover_inv, out, 1, precision);
109  wil_dslash(out, gauge, tmp, 0, dagger, precision, gauge_param);
110  }
111  xpay(in, kappa2, out, Vh*spinorSiteSize, precision);
112  break;
114  wil_dslash(out, gauge, in, 1, dagger, precision, gauge_param);
115  apply_clover(tmp, clover_inv, out, 1, precision);
116  wil_dslash(out, gauge, tmp, 0, dagger, precision, gauge_param);
117  apply_clover(tmp, clover, in, 0, precision);
118  xpay(tmp, kappa2, out, Vh*spinorSiteSize, precision);
119  break;
120  case QUDA_MATPC_ODD_ODD:
121  if (!dagger) {
122  wil_dslash(tmp, gauge, in, 0, dagger, precision, gauge_param);
123  apply_clover(out, clover_inv, tmp, 0, precision);
124  wil_dslash(tmp, gauge, out, 1, dagger, precision, gauge_param);
125  apply_clover(out, clover_inv, tmp, 1, precision);
126  } else {
127  apply_clover(tmp, clover_inv, in, 1, precision);
128  wil_dslash(out, gauge, tmp, 0, dagger, precision, gauge_param);
129  apply_clover(tmp, clover_inv, out, 0, precision);
130  wil_dslash(out, gauge, tmp, 1, dagger, precision, gauge_param);
131  }
132  xpay(in, kappa2, out, Vh*spinorSiteSize, precision);
133  break;
135  wil_dslash(out, gauge, in, 0, dagger, precision, gauge_param);
136  apply_clover(tmp, clover_inv, out, 0, precision);
137  wil_dslash(out, gauge, tmp, 1, dagger, precision, gauge_param);
138  apply_clover(tmp, clover, in, 1, precision);
139  xpay(tmp, kappa2, out, Vh*spinorSiteSize, precision);
140  break;
141  default:
142  errorQuda("Unsupoorted matpc=%d", matpc_type);
143  }
144 
145  free(tmp);
146 }
147 
148 // Apply the full Wilson-clover operator
149 void clover_mat(void *out, void **gauge, void *clover, void *in, double kappa,
150  int dagger, QudaPrecision precision, QudaGaugeParam &gauge_param) {
151 
152  void *tmp = malloc(V*spinorSiteSize*precision);
153 
154  void *inEven = in;
155  void *inOdd = (char*)in + Vh*spinorSiteSize*precision;
156  void *outEven = out;
157  void *outOdd = (char*)out + Vh*spinorSiteSize*precision;
158  void *tmpEven = tmp;
159  void *tmpOdd = (char*)tmp + Vh*spinorSiteSize*precision;
160 
161  // Odd part
162  wil_dslash(outOdd, gauge, inEven, 1, dagger, precision, gauge_param);
163  apply_clover(tmpOdd, clover, inOdd, 1, precision);
164 
165  // Even part
166  wil_dslash(outEven, gauge, inOdd, 0, dagger, precision, gauge_param);
167  apply_clover(tmpEven, clover, inEven, 0, precision);
168 
169  // lastly apply the kappa term
170  xpay(tmp, -kappa, out, V*spinorSiteSize, precision);
171 
172  free(tmp);
173 }
174 
175 void applyTwist(void *out, void *in, void *tmpH, double a, QudaPrecision precision) {
176  switch (precision) {
178  for(int i = 0; i < Vh; i++)
179  for(int s = 0; s < 4; s++) {
180  double a5 = ((s / 2) ? -1.0 : +1.0) * a;
181  for(int c = 0; c < 3; c++) {
182  ((double *) out)[i * 24 + s * 6 + c * 2 + 0] = ((double *) tmpH)[i * 24 + s * 6 + c * 2 + 0] - a5*((double *) in)[i * 24 + s * 6 + c * 2 + 1];
183  ((double *) out)[i * 24 + s * 6 + c * 2 + 1] = ((double *) tmpH)[i * 24 + s * 6 + c * 2 + 1] + a5*((double *) in)[i * 24 + s * 6 + c * 2 + 0];
184  }
185  }
186  break;
188  for(int i = 0; i < Vh; i++)
189  for(int s = 0; s < 4; s++) {
190  float a5 = ((s / 2) ? -1.0 : +1.0) * a;
191  for(int c = 0; c < 3; c++) {
192  ((float *) out)[i * 24 + s * 6 + c * 2 + 0] = ((float *) tmpH)[i * 24 + s * 6 + c * 2 + 0] - a5*((float *) in)[i * 24 + s * 6 + c * 2 + 1];
193  ((float *) out)[i * 24 + s * 6 + c * 2 + 1] = ((float *) tmpH)[i * 24 + s * 6 + c * 2 + 1] + a5*((float *) in)[i * 24 + s * 6 + c * 2 + 0];
194  }
195  }
196  break;
197  default:
198  errorQuda("Unsupported precision %d", precision);
199  }
200 }
201 
202 // Apply (C + i*a*gamma_5)/(C^2 + a^2)
203 void twistCloverGamma5(void *out, void *in, void *clover, void *cInv, const int dagger, const double kappa, const double mu,
204  const QudaTwistFlavorType flavor, const int parity, QudaTwistGamma5Type twist, QudaPrecision precision) {
205  void *tmp1 = malloc(Vh*spinorSiteSize*precision);
206  void *tmp2 = malloc(Vh*spinorSiteSize*precision);
207 
208  double a = 0.0;
209 
210  if (twist == QUDA_TWIST_GAMMA5_DIRECT) {
211  a = 2.0 * kappa * mu * flavor;
212 
213  if (dagger) a *= -1.0;
214 
215  apply_clover(tmp1, clover, in, parity, precision);
216  applyTwist(out, in, tmp1, a, precision);
217  } else if (twist == QUDA_TWIST_GAMMA5_INVERSE) {
218  a = -2.0 * kappa * mu * flavor;
219 
220  if (dagger) a *= -1.0;
221 
222  apply_clover(tmp1, clover, in, parity, precision);
223  applyTwist(tmp2, in, tmp1, a, precision);
224  apply_clover(out, cInv, tmp2, parity, precision);
225  } else {
226  printf("Twist type %d not defined\n", twist);
227  exit(0);
228  }
229 
230  free(tmp2);
231  free(tmp1);
232 }
233 
234 void tmc_dslash(void *out, void **gauge, void *in, void *clover, void *cInv, double kappa, double mu, QudaTwistFlavorType flavor,
236  void *tmp1 = malloc(Vh*spinorSiteSize*precision);
237  void *tmp2 = malloc(Vh*spinorSiteSize*precision);
238 
239  if (dagger) {
240  twistCloverGamma5(tmp1, in, clover, cInv, dagger, kappa, mu, flavor, 1-parity, QUDA_TWIST_GAMMA5_INVERSE, precision);
241  if (matpc_type == QUDA_MATPC_EVEN_EVEN_ASYMMETRIC || matpc_type == QUDA_MATPC_ODD_ODD_ASYMMETRIC) {
242  wil_dslash(tmp2, gauge, tmp1, parity, dagger, precision, param);
243  twistCloverGamma5(out, tmp2, clover, cInv, dagger, kappa, mu, flavor, parity, QUDA_TWIST_GAMMA5_INVERSE, precision);
244  } else {
245  wil_dslash(out, gauge, tmp1, parity, dagger, precision, param);
246  }
247  } else {
248  wil_dslash(tmp1, gauge, in, parity, dagger, precision, param);
249  twistCloverGamma5(out, tmp1, clover, cInv, dagger, kappa, mu, flavor, parity, QUDA_TWIST_GAMMA5_INVERSE, precision);
250  }
251 
252  free(tmp2);
253  free(tmp1);
254 }
255 
256 // Apply the full twisted-clover operator
257 void tmc_mat(void *out, void **gauge, void *clover, void *in, double kappa, double mu,
259 
260  void *tmp = malloc(V*spinorSiteSize*precision);
261 
262  void *inEven = in;
263  void *inOdd = (char*)in + Vh*spinorSiteSize*precision;
264  void *outEven = out;
265  void *outOdd = (char*)out + Vh*spinorSiteSize*precision;
266  void *tmpEven = tmp;
267  void *tmpOdd = (char*)tmp + Vh*spinorSiteSize*precision;
268 
269  // Odd part
270  wil_dslash(outOdd, gauge, inEven, 1, dagger, precision, gauge_param);
271  twistCloverGamma5(tmpOdd, inOdd, clover, NULL, dagger, kappa, mu, flavor, 1, QUDA_TWIST_GAMMA5_DIRECT, precision);
272 
273  // Even part
274  wil_dslash(outEven, gauge, inOdd, 0, dagger, precision, gauge_param);
275  twistCloverGamma5(tmpEven, inEven, clover, NULL, dagger, kappa, mu, flavor, 0, QUDA_TWIST_GAMMA5_DIRECT, precision);
276 
277  // lastly apply the kappa term
278  xpay(tmp, -kappa, out, V*spinorSiteSize, precision);
279 
280  free(tmp);
281 }
282 
283 // Apply the even-odd preconditioned Dirac operator
284 void tmc_matpc(void *out, void **gauge, void *in, void *clover, void *cInv, double kappa, double mu, QudaTwistFlavorType flavor,
286 
287  double kappa2 = -kappa*kappa;
288 
289  void *tmp1 = malloc(Vh*spinorSiteSize*precision);
290  void *tmp2 = malloc(Vh*spinorSiteSize*precision);
291 
292  switch(matpc_type) {
294  if (!dagger) {
295  wil_dslash(out, gauge, in, 1, dagger, precision, gauge_param);
296  twistCloverGamma5(tmp1, out, clover, cInv, dagger, kappa, mu, flavor, 1, QUDA_TWIST_GAMMA5_INVERSE, precision);
297  wil_dslash(tmp2, gauge, tmp1, 0, dagger, precision, gauge_param);
298  twistCloverGamma5(out, tmp2, clover, cInv, dagger, kappa, mu, flavor, 0, QUDA_TWIST_GAMMA5_INVERSE, precision);
299  } else {
300  twistCloverGamma5(out, in, clover, cInv, dagger, kappa, mu, flavor, 0, QUDA_TWIST_GAMMA5_INVERSE, precision);
301  wil_dslash(tmp1, gauge, out, 1, dagger, precision, gauge_param);
302  twistCloverGamma5(tmp2, tmp1, clover, cInv, dagger, kappa, mu, flavor, 1, QUDA_TWIST_GAMMA5_INVERSE, precision);
303  wil_dslash(out, gauge, tmp2, 0, dagger, precision, gauge_param);
304  }
305  xpay(in, kappa2, out, Vh*spinorSiteSize, precision);
306  break;
308  wil_dslash(tmp1, gauge, in, 1, dagger, precision, gauge_param);
309  twistCloverGamma5(tmp2, tmp1, clover, cInv, dagger, kappa, mu, flavor, 1, QUDA_TWIST_GAMMA5_INVERSE, precision);
310  wil_dslash(out, gauge, tmp2, 0, dagger, precision, gauge_param);
311  twistCloverGamma5(tmp2, in, clover, cInv, dagger, kappa, mu, flavor, 0, QUDA_TWIST_GAMMA5_DIRECT, precision);
312  xpay(tmp2, kappa2, out, Vh*spinorSiteSize, precision);
313  break;
314  case QUDA_MATPC_ODD_ODD:
315  if (!dagger) {
316  wil_dslash(out, gauge, in, 0, dagger, precision, gauge_param);
317  twistCloverGamma5(tmp1, out, clover, cInv, dagger, kappa, mu, flavor, 0, QUDA_TWIST_GAMMA5_INVERSE, precision);
318  wil_dslash(tmp2, gauge, tmp1, 1, dagger, precision, gauge_param);
319  twistCloverGamma5(out, tmp2, clover, cInv, dagger, kappa, mu, flavor, 1, QUDA_TWIST_GAMMA5_INVERSE, precision);
320  } else {
321  twistCloverGamma5(out, in, clover, cInv, dagger, kappa, mu, flavor, 1, QUDA_TWIST_GAMMA5_INVERSE, precision);
322  wil_dslash(tmp1, gauge, out, 0, dagger, precision, gauge_param);
323  twistCloverGamma5(tmp2, tmp1, clover, cInv, dagger, kappa, mu, flavor, 0, QUDA_TWIST_GAMMA5_INVERSE, precision);
324  wil_dslash(out, gauge, tmp2, 1, dagger, precision, gauge_param);
325  }
326  xpay(in, kappa2, out, Vh*spinorSiteSize, precision);
327  break;
329  wil_dslash(tmp1, gauge, in, 0, dagger, precision, gauge_param);
330  twistCloverGamma5(tmp2, tmp1, clover, cInv, dagger, kappa, mu, flavor, 0, QUDA_TWIST_GAMMA5_INVERSE, precision);
331  wil_dslash(out, gauge, tmp2, 1, dagger, precision, gauge_param);
332  twistCloverGamma5(tmp1, in, clover, cInv, dagger, kappa, mu, flavor, 1, QUDA_TWIST_GAMMA5_DIRECT, precision);
333  xpay(tmp1, kappa2, out, Vh*spinorSiteSize, precision);
334  break;
335  default:
336  errorQuda("Unsupported matpc=%d", matpc_type);
337  }
338 
339  free(tmp2);
340  free(tmp1);
341 }
cudaColorSpinorField * tmp2
cudaColorSpinorField * tmp1
double mu
Definition: test_util.cpp:1648
enum QudaPrecision_s QudaPrecision
double kappa
Definition: test_util.cpp:1647
Matrix< N, std::complex< T > > conj(const Matrix< N, std::complex< T > > &mat)
#define errorQuda(...)
Definition: util_quda.h:121
void tmc_dslash(void *out, void **gauge, void *in, void *clover, void *cInv, double kappa, double mu, QudaTwistFlavorType flavor, int parity, QudaMatPCType matpc_type, int dagger, QudaPrecision precision, QudaGaugeParam &param)
cudaColorSpinorField * tmp
Definition: covdev_test.cpp:44
void clover_matpc(void *out, void **gauge, void *clover, void *clover_inv, void *in, double kappa, QudaMatPCType matpc_type, int dagger, QudaPrecision precision, QudaGaugeParam &gauge_param)
QudaGaugeParam gauge_param
void tmc_mat(void *out, void **gauge, void *clover, void *in, double kappa, double mu, QudaTwistFlavorType flavor, int dagger, QudaPrecision precision, QudaGaugeParam &gauge_param)
#define spinorSiteSize
void apply_clover(void *out, void *clover, void *in, int parity, QudaPrecision precision)
void xpay(ColorSpinorField &x, double a, ColorSpinorField &y)
Definition: blas_quda.h:37
QudaGaugeParam param
Definition: pack_test.cpp:17
const int nColor
Definition: covdev_test.cpp:75
cpuColorSpinorField * in
enum QudaMatPCType_s QudaMatPCType
void clover_mat(void *out, void **gauge, void *clover, void *in, double kappa, int dagger, QudaPrecision precision, QudaGaugeParam &gauge_param)
QudaMatPCType matpc_type
Definition: test_util.cpp:1662
void applyTwist(void *out, void *in, void *tmpH, double a, QudaPrecision precision)
int V
Definition: test_util.cpp:27
void cloverReference(sFloat *out, cFloat *clover, sFloat *in, int parity)
Apply the clover matrix field.
void clover_dslash(void *out, void **gauge, void *clover, void *in, int parity, int dagger, QudaPrecision precision, QudaGaugeParam &param)
cpuColorSpinorField * out
__shared__ float s[]
void tmc_matpc(void *out, void **gauge, void *in, void *clover, void *cInv, double kappa, double mu, QudaTwistFlavorType flavor, QudaMatPCType matpc_type, int dagger, QudaPrecision precision, QudaGaugeParam &gauge_param)
void twistCloverGamma5(void *out, void *in, void *clover, void *cInv, const int dagger, const double kappa, const double mu, const QudaTwistFlavorType flavor, const int parity, QudaTwistGamma5Type twist, QudaPrecision precision)
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
void wil_dslash(void *out, void **gauge, void *in, int oddBit, int daggerBit, QudaPrecision precision, QudaGaugeParam &gauge_param)
QudaDagType dagger
Definition: test_util.cpp:1620
QudaParity parity
Definition: covdev_test.cpp:54
int Vh
Definition: test_util.cpp:28
enum QudaTwistFlavorType_s QudaTwistFlavorType