QUDA  v0.7.0
A library for QCD on GPUs
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
dirac_twisted_clover.cpp
Go to the documentation of this file.
1 #include <dirac_quda.h>
2 #include <blas_quda.h>
3 #include <iostream>
4 
5 namespace quda {
6 
7  namespace twistedclover {
8 #include <dslash_init.cuh>
9  }
10 
11  namespace dslash_aux {
12 #include <dslash_init.cuh>
13  }
14 
16  : DiracWilson(param, nDim), mu(param.mu), epsilon(param.epsilon), clover(*(param.clover)), cloverInv(*(param.cloverInv))
17  {
18  twistedclover::initConstants(*param.gauge,profile);
19  dslash_aux::initConstants(*param.gauge,profile);
20  }
21 
23  : DiracWilson(dirac), mu(dirac.mu), epsilon(dirac.epsilon), clover(dirac.clover), cloverInv(dirac.cloverInv)
24  {
25  twistedclover::initConstants(dirac.gauge,profile);
26  dslash_aux::initConstants(dirac.gauge,profile);
27  }
28 
30 
32  {
33  if (&dirac != this)
34  {
36  clover = dirac.clover;
37  cloverInv = dirac.cloverInv;
38  }
39 
40  return *this;
41  }
42 
44  {
45  Dirac::checkParitySpinor(out, in);
46 
47  if (out.Volume() != clover.VolumeCB())
48  errorQuda("Parity spinor volume %d doesn't match clover checkboard volume %d", out.Volume(), clover.VolumeCB());
49  }
50 
51  // Protected method for applying twist
52 
54  {
55  checkParitySpinor(out, in);
56 
58  errorQuda("Twist flavor not set %d\n", in.TwistFlavor());
59 
61  {
62  FullClover *cs = new FullClover(clover);
63  FullClover *cI = new FullClover(cloverInv, false);
64  double flavor_mu = in.TwistFlavor() * mu;
65  twistCloverGamma5Cuda(&out, &in, dagger, kappa, flavor_mu, 0.0, twistType, cs, cI, parity);
66 
67  if (twistType == QUDA_TWIST_GAMMA5_INVERSE)
68  flops += 1056ll*in.Volume();
69  else
70  flops += 552ll*in.Volume();
71 
72  delete cs;
73  delete cI;
74  }
75  else
76  errorQuda("DiracTwistedClover::twistedCloverApply method for flavor doublet is not implemented..\n");
77  }
78 
79 
80  // Public method to apply the twist
82  {
84  }
85 
87  {
88  checkFullSpinor(out, in);
89  if (in.TwistFlavor() != out.TwistFlavor())
90  errorQuda("Twist flavors %d %d don't match", in.TwistFlavor(), out.TwistFlavor());
91 
93  errorQuda("Twist flavor not set %d\n", in.TwistFlavor());
94  }
95 
96  // We can eliminate this temporary at the expense of more kernels (like clover)
97  cudaColorSpinorField *tmp=0; // this hack allows for tmp2 to be full or parity field
98  if (tmp2) {
99  if (tmp2->SiteSubset() == QUDA_FULL_SITE_SUBSET) tmp = &(tmp2->Even());
100  else tmp = tmp2;
101  }
102  bool reset = newTmp(&tmp, in.Even());
103 
104  twistedclover::setFace(face1,face2); // FIXME: temporary hack maintain C linkage for dslashCuda
105 
106  FullClover *cs = new FullClover(clover);
107  FullClover *cI = new FullClover(cloverInv, false);
108 
110  double a = 2.0 * kappa * in.TwistFlavor() * mu;//for direct twist (must be daggered separately)
113  flops += (1320ll+552ll)*in.Volume();
114  } else {
115  errorQuda("Non-deg twisted clover not implemented yet");
116  }
117  deleteTmp(&tmp, reset);
118  delete cs;
119  delete cI;
120  }
121 
123  {
124  checkFullSpinor(out, in);
125  bool reset = newTmp(&tmp1, in);
126 
127  M(*tmp1, in);
128  Mdag(out, *tmp1);
129 
130  deleteTmp(&tmp1, reset);
131  }
132 
135  const QudaSolutionType solType) const
136  {
137  if (solType == QUDA_MATPC_SOLUTION || solType == QUDA_MATPCDAG_MATPC_SOLUTION) {
138  errorQuda("Preconditioned solution requires a preconditioned solve_type");
139  }
140 
141  src = &b;
142  sol = &x;
143  }
144 
146  const QudaSolutionType solType) const
147  {
148  // do nothing
149  }
150 
151 
153 
155 
157  {
158 
159  }
160 
162  {
163  if (&dirac != this) {
165  }
166  return *this;
167  }
168 
169  // Public method to apply the inverse twist
171  {
173  }
174 
175  // apply hopping term, then inverse twist: (A_ee^-1 D_eo) or (A_oo^-1 D_oe),
176  // and likewise for dagger: (D^dagger_eo D_ee^-1) or (D^dagger_oe A_oo^-1)
179  {
180  checkParitySpinor(in, out);
181  checkSpinorAlias(in, out);
182 
183  if (in.TwistFlavor() != out.TwistFlavor())
184  errorQuda("Twist flavors %d %d don't match", in.TwistFlavor(), out.TwistFlavor());
186  errorQuda("Twist flavor not set %d\n", in.TwistFlavor());
187 
188  twistedclover::setFace(face1,face2); // FIXME: temporary hack maintain C linkage for dslashCuda
189 
190  FullClover *cs = new FullClover(clover);
191  FullClover *cI = new FullClover(cloverInv, false);
192 
194  double a = -2.0 * kappa * in.TwistFlavor() * mu; //for invert twist (not daggered)
195  double b = 1.;// / (1.0 + a*a); //for invert twist
196  if (!dagger || matpcType == QUDA_MATPC_EVEN_EVEN_ASYMMETRIC || matpcType == QUDA_MATPC_ODD_ODD_ASYMMETRIC) {
197  twistedCloverDslashCuda(&out, gauge, cs, cI, &in, parity, dagger, 0, QUDA_DEG_DSLASH_CLOVER_TWIST_INV, a, b, 0.0, 0.0, commDim, profile);
198  flops += 2376ll*in.Volume();
199  } else {
200  twistedCloverDslashCuda(&out, gauge, cs, cI, &in, parity, dagger, 0, QUDA_DEG_CLOVER_TWIST_INV_DSLASH, a, b, 0.0, 0.0, commDim, profile);
201  flops += 1320ll*in.Volume();
202  }
203  } else {//TWIST doublet :
204  errorQuda("Non-degenerate DiracTwistedCloverPC is not implemented \n");
205  }
206  delete cs;
207  delete cI;
208  }
209 
210  // xpay version of the above
212  (cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity, const cudaColorSpinorField &x, const double &k) const
213  {
214  checkParitySpinor(in, out);
215  checkSpinorAlias(in, out);
216  if (in.TwistFlavor() != out.TwistFlavor())
217  errorQuda("Twist flavors %d %d don't match", in.TwistFlavor(), out.TwistFlavor());
219  errorQuda("Twist flavor not set %d\n", in.TwistFlavor());
220 
221  twistedclover::setFace(face1,face2); // FIXME: temporary hack maintain C linkage for dslashCuda
222 
223  FullClover *cs = new FullClover(clover);
224  FullClover *cI = new FullClover(cloverInv, false);
225 
227  double a = -2.0 * kappa * in.TwistFlavor() * mu; //for invert twist
228  // double b = k / (1.0 + a*a); //for invert twist NO HABRÍA QUE APLICAR CLOVER_TWIST_INV???
229  double b = k; //for invert twist NO HABRÍA QUE APLICAR CLOVER_TWIST_INV???
230  if (!dagger) {
231  twistedCloverDslashCuda(&out, gauge, cs, cI, &in, parity, dagger, &x, QUDA_DEG_DSLASH_CLOVER_TWIST_INV, a, b, 0.0, 0.0, commDim, profile);
232  flops += 2400ll*in.Volume();
233  } else { // tmp1 can alias in, but tmp2 can alias x so must not use this
234  twistedCloverDslashCuda(&out, gauge, cs, cI, &in, parity, dagger, &x, QUDA_DEG_CLOVER_TWIST_INV_DSLASH, a, b, 0.0, 0.0, commDim, profile);
235  flops += 1344ll*in.Volume();
236  }
237  } else {//TWIST_DOUBLET:
238  errorQuda("Non-degenerate DiracTwistedCloverPC is not implemented \n");
239  }
240  delete cs;
241  delete cI;
242  }
243 
245  {
246  double kappa2 = -kappa*kappa;
247 
248  bool reset = newTmp(&tmp1, in);
249 
250  FullClover *cs = new FullClover(clover);
251  FullClover *cI = new FullClover(cloverInv, false);
252 
255  if (dagger) {
257  Dslash(out, *tmp1, QUDA_ODD_PARITY);
259  DslashXpay(out, *tmp1, QUDA_EVEN_PARITY, in, kappa2);
260  } else {
261  Dslash(*tmp1, in, QUDA_ODD_PARITY);
262  DslashXpay(out, *tmp1, QUDA_EVEN_PARITY, in, kappa2);
263  }
264  } else if (matpcType == QUDA_MATPC_ODD_ODD) {
265  if (dagger) {
267  Dslash(out, *tmp1, QUDA_EVEN_PARITY);
269  DslashXpay(out, *tmp1, QUDA_ODD_PARITY, in, kappa2);
270  } else {
272  DslashXpay(out, *tmp1, QUDA_ODD_PARITY, in, kappa2);
273  }
274  } else {//asymmetric preconditioning
275  double a = 2.0 * kappa * in.TwistFlavor() * mu;
277  Dslash(*tmp1, in, QUDA_ODD_PARITY);
279  flops += (1320ll+96ll)*in.Volume();
283  flops += (1320ll+96ll)*in.Volume();
284  }else { // symmetric preconditioning
285  errorQuda("Invalid matpcType");
286  }
287  }
288  } else { //Twist doublet
289  errorQuda("Non-degenerate DiracTwistedCloverPC is not implemented \n");
290  }
291 
292  delete cs;
293  delete cI;
294 
295  deleteTmp(&tmp1, reset);
296  }
297 
299  {
300  // need extra temporary because of symmetric preconditioning dagger
301  bool reset = newTmp(&tmp2, in);
302  M(*tmp2, in);
303  Mdag(out, *tmp2);
304  deleteTmp(&tmp2, reset);
305  }
306 
309  const QudaSolutionType solType) const
310  {
311  // we desire solution to preconditioned system
312  if (solType == QUDA_MATPC_SOLUTION || solType == QUDA_MATPCDAG_MATPC_SOLUTION) {
313  src = &b;
314  sol = &x;
315  return;
316  }
317 
318  bool reset = newTmp(&tmp1, b.Even());
319 
320  // we desire solution to full system
323  // src = A_ee^-1 (b_e + k D_eo A_oo^-1 b_o)
324  src = &(x.Odd());
325  TwistCloverInv(*src, b.Odd(), 1);
327  TwistCloverInv(*src, *tmp1, 0);
328  sol = &(x.Even());
329  } else if (matpcType == QUDA_MATPC_ODD_ODD) {
330  // src = A_oo^-1 (b_o + k D_oe A_ee^-1 b_e)
331  src = &(x.Even());
332  TwistCloverInv(*src, b.Even(), 0);
334  TwistCloverInv(*src, *tmp1, 1);
335  sol = &(x.Odd());
337  // src = b_e + k D_eo A_oo^-1 b_o
338  src = &(x.Odd());
339  TwistCloverInv(*tmp1, b.Odd(), 1); // safe even when *tmp1 = b.odd
341  sol = &(x.Even());
343  // src = b_o + k D_oe A_ee^-1 b_e
344  src = &(x.Even());
345  TwistCloverInv(*tmp1, b.Even(), 0); // safe even when *tmp1 = b.even
347  sol = &(x.Odd());
348  } else {
349  errorQuda("MatPCType %d not valid for DiracTwistedCloverPC", matpcType);
350  }
351  } else {//doublet:
352  errorQuda("Non-degenrate DiracTwistedCloverPC is not implemented \n");
353  }//end of doublet
354  // here we use final solution to store parity solution and parity source
355  // b is now up for grabs if we want
356 
357  deleteTmp(&tmp1, reset);
358  }
359 
361  const QudaSolutionType solType) const
362  {
363  if (solType == QUDA_MATPC_SOLUTION || solType == QUDA_MATPCDAG_MATPC_SOLUTION) {
364  return;
365  }
366 
367  checkFullSpinor(x, b);
368  bool reset = newTmp(&tmp1, b.Even());
369 
370  // create full solution
373  // x_o = A_oo^-1 (b_o + k D_oe x_e)
375  TwistCloverInv(x.Odd(), *tmp1, 1);
377  // x_e = A_ee^-1 (b_e + k D_eo x_o)
379  TwistCloverInv(x.Even(), *tmp1, 0);
380  } else {
381  errorQuda("MatPCType %d not valid for DiracTwistedCloverPC", matpcType);
382  }
383  } else { //twist doublet:
384  errorQuda("Non-degenrate DiracTwistedCloverPC is not implemented \n");
385  }//end of twist doublet...
386  deleteTmp(&tmp1, reset);
387  }
388 } // namespace quda
FaceBuffer face1
Definition: dirac_quda.h:148
int commDim(int)
virtual void reconstruct(cudaColorSpinorField &x, const cudaColorSpinorField &b, const QudaSolutionType) const
cudaGaugeField & gauge
Definition: dirac_quda.h:88
virtual void checkParitySpinor(const cudaColorSpinorField &, const cudaColorSpinorField &) const
Definition: dirac.cpp:84
unsigned long long flops
Definition: dirac_quda.h:93
void reconstruct(cudaColorSpinorField &x, const cudaColorSpinorField &b, const QudaSolutionType) const
int VolumeCB() const
#define errorQuda(...)
Definition: util_quda.h:73
void twistedCloverApply(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaTwistGamma5Type twistType, const int parity) const
__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int mu
bool newTmp(cudaColorSpinorField **, const cudaColorSpinorField &) const
Definition: dirac.cpp:51
TimeProfile profile
Definition: dirac_quda.h:104
DiracWilson & operator=(const DiracWilson &dirac)
virtual void M(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
DiracTwistedCloverPC(const DiracTwistedCloverPC &dirac)
QudaDagType dagger
Definition: test_util.cpp:1558
cudaGaugeField * gauge
Definition: dirac_quda.h:30
QudaGaugeParam param
Definition: pack_test.cpp:17
cudaColorSpinorField & Odd() const
int commDim[QUDA_MAX_DIM]
Definition: dirac_quda.h:102
void TwistClover(cudaColorSpinorField &out, const cudaColorSpinorField &in, const int parity) const
cudaColorSpinorField * tmp
void M(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
VOLATILE spinorFloat kappa
cudaCloverField & clover
Definition: dirac_quda.h:420
DiracTwistedCloverPC & operator=(const DiracTwistedCloverPC &dirac)
virtual void MdagM(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
cpuColorSpinorField * in
DiracTwistedClover & operator=(const DiracTwistedClover &dirac)
enum QudaSolutionType_s QudaSolutionType
void twistedCloverDslashCuda(cudaColorSpinorField *out, const cudaGaugeField &gauge, const FullClover *clover, const FullClover *cloverInv, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const QudaTwistCloverDslashType type, const double &kappa, const double &mu, const double &epsilon, const double &k, const int *commDim, TimeProfile &profile, const QudaDslashPolicy &dslashPolicy=QUDA_DSLASH2)
void deleteTmp(cudaColorSpinorField **, const bool &reset) const
Definition: dirac.cpp:59
DiracTwistedClover(const DiracTwistedClover &dirac)
Dirac * dirac
Definition: dslash_test.cpp:45
void prepare(cudaColorSpinorField *&src, cudaColorSpinorField *&sol, cudaColorSpinorField &x, cudaColorSpinorField &b, const QudaSolutionType) const
QudaDagType dagger
Definition: dirac_quda.h:92
void MdagM(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
enum QudaParity_s QudaParity
virtual void DslashXpay(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity, const cudaColorSpinorField &x, const double &k) const
int x[4]
double kappa
Definition: dirac_quda.h:89
QudaMatPCType matpcType
Definition: dirac_quda.h:91
virtual void prepare(cudaColorSpinorField *&src, cudaColorSpinorField *&sol, cudaColorSpinorField &x, cudaColorSpinorField &b, const QudaSolutionType) const
FaceBuffer face2
Definition: dirac_quda.h:148
cudaColorSpinorField * tmp2
Definition: dirac_quda.h:95
virtual void checkFullSpinor(const cudaColorSpinorField &, const cudaColorSpinorField &) const
Definition: dirac.cpp:121
virtual void DslashXpay(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity, const cudaColorSpinorField &x, const double &k) const
void twistCloverGamma5Cuda(cudaColorSpinorField *out, const cudaColorSpinorField *in, const int dagger, const double &kappa, const double &mu, const double &epsilon, const QudaTwistGamma5Type twist, const FullClover *clov, const FullClover *clovInv, const int parity)
Definition: dslash_quda.cu:495
cpuColorSpinorField * out
void TwistCloverInv(cudaColorSpinorField &out, const cudaColorSpinorField &in, const int parity) const
cudaColorSpinorField * tmp1
Definition: dirac_quda.h:94
QudaTwistFlavorType TwistFlavor() const
cudaCloverField & cloverInv
Definition: dirac_quda.h:421
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
virtual void Dslash(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity) const
void Mdag(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
Definition: dirac.cpp:68
QudaSiteSubset SiteSubset() const
const QudaParity parity
Definition: dslash_test.cpp:29
void * gauge[4]
Definition: su3_test.cpp:15
cudaColorSpinorField & Even() const
void checkParitySpinor(const cudaColorSpinorField &, const cudaColorSpinorField &) const