QUDA  1.0.0
•All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
dirac_quda.h
Go to the documentation of this file.
1 #ifndef _DIRAC_QUDA_H
2 #define _DIRAC_QUDA_H
3 
4 #include <quda_internal.h>
5 #include <color_spinor_field.h>
6 #include <gauge_field.h>
7 #include <clover_field.h>
8 #include <dslash_quda.h>
9 #include <blas_quda.h>
10 
11 #include <typeinfo>
12 
13 namespace quda {
14 
15  class Transfer;
16  class Dirac;
17 
18  // Params for Dirac operator
19  class DiracParam {
20 
21  public:
22  QudaDiracType type;
23  double kappa;
24  double mass;
25  double m5; // used by domain wall only
26  int Ls; // used by domain wall and twisted mass
27  Complex b_5[QUDA_MAX_DWF_LS]; // used by mobius domain wall only
28  Complex c_5[QUDA_MAX_DWF_LS]; // used by mobius domain wall only
29  QudaMatPCType matpcType;
30  QudaDagType dagger;
31  cudaGaugeField *gauge;
32  cudaGaugeField *fatGauge; // used by staggered only
33  cudaGaugeField *longGauge; // used by staggered only
34  int laplace3D;
35  cudaCloverField *clover;
36 
37  double mu; // used by twisted mass only
38  double mu_factor; // used by multigrid only
39  double epsilon; //2nd tm parameter (used by twisted mass only)
40 
41  ColorSpinorField *tmp1;
42  ColorSpinorField *tmp2; // used by Wilson-like kernels only
43 
44  int commDim[QUDA_MAX_DIM]; // whether to do comms or not
45 
46  QudaPrecision halo_precision; // only does something for DiracCoarse at present
47 
48  // for multigrid only
49  Transfer *transfer;
50  Dirac *dirac;
51  bool need_bidirectional; // whether or not we need to force a bi-directional build
52 
53  DiracParam() :
54  type(QUDA_INVALID_DIRAC),
55  kappa(0.0),
56  m5(0.0),
57  matpcType(QUDA_MATPC_INVALID),
58  dagger(QUDA_DAG_INVALID),
59  gauge(0),
60  clover(0),
61  mu(0.0),
62  mu_factor(0.0),
63  epsilon(0.0),
64  tmp1(0),
65  tmp2(0),
66  halo_precision(QUDA_INVALID_PRECISION),
67  need_bidirectional(false)
68  {
69  for (int i=0; i<QUDA_MAX_DIM; i++) commDim[i] = 1;
70  }
71 
72  void print() {
73  printfQuda("Printing DslashParam\n");
74  printfQuda("type = %d\n", type);
75  printfQuda("kappa = %g\n", kappa);
76  printfQuda("mass = %g\n", mass);
77  printfQuda("laplace3D = %d\n", laplace3D);
78  printfQuda("m5 = %g\n", m5);
79  printfQuda("Ls = %d\n", Ls);
80  printfQuda("matpcType = %d\n", matpcType);
81  printfQuda("dagger = %d\n", dagger);
82  printfQuda("mu = %g\n", mu);
83  printfQuda("epsilon = %g\n", epsilon);
84  printfQuda("halo_precision = %d\n", halo_precision);
85  for (int i=0; i<QUDA_MAX_DIM; i++) printfQuda("commDim[%d] = %d\n", i, commDim[i]);
86  for (int i = 0; i < Ls; i++)
87  printfQuda(
88  "b_5[%d] = %e %e \t c_5[%d] = %e %e\n", i, b_5[i].real(), b_5[i].imag(), i, c_5[i].real(), c_5[i].imag());
89  }
90 
91  };
92 
93  void setDiracParam(DiracParam &diracParam, QudaInvertParam *inv_param, bool pc);
94  void setDiracSloppyParam(DiracParam &diracParam, QudaInvertParam *inv_param, bool pc);
95 
96  // forward declarations
97  class DiracMatrix;
98  class DiracM;
99  class DiracMdagM;
100  class DiracMMdag;
101  class DiracMdag;
102  //Forward declaration of multigrid Transfer class
103  class Transfer;
104 
105  // Abstract base class
106  class Dirac : public Object {
107 
108  friend class DiracMatrix;
109  friend class DiracM;
110  friend class DiracMdagM;
111  friend class DiracMMdag;
112  friend class DiracMdag;
113 
114  protected:
115  cudaGaugeField *gauge;
116  double kappa;
117  double mass;
118  int laplace3D;
119  QudaMatPCType matpcType;
120  mutable QudaDagType dagger; // mutable to simplify implementation of Mdag
121  mutable unsigned long long flops;
122  mutable ColorSpinorField *tmp1; // temporary hack
123  mutable ColorSpinorField *tmp2; // temporary hack
124  QudaDiracType type;
125  mutable QudaPrecision halo_precision; // only does something for DiracCoarse at present
126 
127  bool newTmp(ColorSpinorField **, const ColorSpinorField &) const;
128  void deleteTmp(ColorSpinorField **, const bool &reset) const;
129 
130  mutable int commDim[QUDA_MAX_DIM]; // whether do comms or not
131 
132  mutable TimeProfile profile;
133 
134  public:
135  Dirac(const DiracParam &param);
136  Dirac(const Dirac &dirac);
137  virtual ~Dirac();
138  Dirac& operator=(const Dirac &dirac);
139 
145  void setCommDim(const int commDim_[QUDA_MAX_DIM]) const {
146  for (int i=0; i<QUDA_MAX_DIM; i++) { commDim[i] = commDim_[i]; }
147  }
148 
149  virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
150  virtual void checkFullSpinor(const ColorSpinorField &, const ColorSpinorField &) const;
151  void checkSpinorAlias(const ColorSpinorField &, const ColorSpinorField &) const;
152 
153  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
154  const QudaParity parity) const = 0;
155  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
156  const QudaParity parity, const ColorSpinorField &x,
157  const double &k) const = 0;
158  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const = 0;
159  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const = 0;
160  void Mdag(ColorSpinorField &out, const ColorSpinorField &in) const;
161  void MMdag(ColorSpinorField &out, const ColorSpinorField &in) const;
162 
163  // required methods to use e-o preconditioning for solving full system
164  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
165  ColorSpinorField &x, ColorSpinorField &b,
166  const QudaSolutionType) const = 0;
167  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
168  const QudaSolutionType) const = 0;
169  void setMass(double mass){ this->mass = mass;}
170  // Dirac operator factory
171  static Dirac* create(const DiracParam &param);
172 
173  double Kappa() const { return kappa; }
174  virtual double Mu() const { return 0.; }
175  virtual double MuFactor() const { return 0.; }
176 
177  unsigned long long Flops() const { unsigned long long rtn = flops; flops = 0; return rtn; }
178 
179 
180  QudaMatPCType getMatPCType() const { return matpcType; }
181  int getStencilSteps() const;
182  void Dagger(QudaDagType dag) const { dagger = dag; }
183  void flipDagger() const { dagger = (dagger == QUDA_DAG_YES) ? QUDA_DAG_NO : QUDA_DAG_YES; }
184 
196  virtual void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
197  double kappa, double mass=0., double mu=0., double mu_factor=0.) const
198  {errorQuda("Not implemented");}
199 
200  QudaPrecision HaloPrecision() const { return halo_precision; }
201  void setHaloPrecision(QudaPrecision halo_precision_) const { halo_precision = halo_precision_; }
202  };
203 
204  // Full Wilson
205  class DiracWilson : public Dirac {
206 
207  protected:
208  void initConstants();
209 
210  public:
211  DiracWilson(const DiracParam &param);
212  DiracWilson(const DiracWilson &dirac);
213  DiracWilson(const DiracParam &param, const int nDims);//to correctly adjust face for DW and non-deg twisted mass
214 
215  virtual ~DiracWilson();
216  DiracWilson& operator=(const DiracWilson &dirac);
217 
218  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
219  const QudaParity parity) const;
220  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
221  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
222  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
223  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
224 
225  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
226  ColorSpinorField &x, ColorSpinorField &b,
227  const QudaSolutionType) const;
228  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
229  const QudaSolutionType) const;
230 
240  virtual void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
241  double kappa, double mass=0.,double mu=0., double mu_factor=0.) const;
242  };
243 
244  // Even-odd preconditioned Wilson
245  class DiracWilsonPC : public DiracWilson {
246 
247  private:
248 
249  public:
250  DiracWilsonPC(const DiracParam &param);
251  DiracWilsonPC(const DiracWilsonPC &dirac);
252  virtual ~DiracWilsonPC();
253  DiracWilsonPC& operator=(const DiracWilsonPC &dirac);
254 
255  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
256  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
257 
258  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
259  ColorSpinorField &x, ColorSpinorField &b,
260  const QudaSolutionType) const;
261  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
262  const QudaSolutionType) const;
263  };
264 
265  // Full clover
266  class DiracClover : public DiracWilson {
267 
268  protected:
269  cudaCloverField &clover;
270  void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
271  void initConstants();
272 
273  public:
274  DiracClover(const DiracParam &param);
275  DiracClover(const DiracClover &dirac);
276  virtual ~DiracClover();
277  DiracClover& operator=(const DiracClover &dirac);
278 
279  void Clover(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
280  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity,
281  const ColorSpinorField &x, const double &k) const;
282  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
283  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
284 
285  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
286  ColorSpinorField &x, ColorSpinorField &b,
287  const QudaSolutionType) const;
288  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
289  const QudaSolutionType) const;
290 
300  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
301  double kappa, double mass=0., double mu=0., double mu_factor=0.) const;
302  };
303 
304  // Even-odd preconditioned clover
305  class DiracCloverPC : public DiracClover {
306 
307  public:
308  DiracCloverPC(const DiracParam &param);
309  DiracCloverPC(const DiracCloverPC &dirac);
310  virtual ~DiracCloverPC();
311  DiracCloverPC& operator=(const DiracCloverPC &dirac);
312 
313  void CloverInv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
314  void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
315  const QudaParity parity) const;
316  void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
317  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
318 
319  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
320  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
321 
322  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
323  ColorSpinorField &x, ColorSpinorField &b,
324  const QudaSolutionType) const;
325  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
326  const QudaSolutionType) const;
327 
340  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
341  double kappa, double mass=0., double mu=0., double mu_factor=0.) const;
342  };
343 
344 
345  // Full domain wall
346  class DiracDomainWall : public DiracWilson {
347 
348  protected:
349  double m5;
350  double kappa5;
351  int Ls; // length of the fifth dimension
352  void checkDWF(const ColorSpinorField &out, const ColorSpinorField &in) const;
353 
354 public:
355  DiracDomainWall(const DiracParam &param);
356  DiracDomainWall(const DiracDomainWall &dirac);
357  virtual ~DiracDomainWall();
358  DiracDomainWall& operator=(const DiracDomainWall &dirac);
359 
360  void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
361  const QudaParity parity) const;
362  void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
363  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
364 
365  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
366  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
367 
368  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
369  ColorSpinorField &x, ColorSpinorField &b,
370  const QudaSolutionType) const;
371  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
372  const QudaSolutionType) const;
373  };
374 
375  // 5d Even-odd preconditioned domain wall
376  class DiracDomainWallPC : public DiracDomainWall {
377 
378  private:
379 
380  public:
381  DiracDomainWallPC(const DiracParam &param);
382  DiracDomainWallPC(const DiracDomainWallPC &dirac);
383  virtual ~DiracDomainWallPC();
384  DiracDomainWallPC& operator=(const DiracDomainWallPC &dirac);
385 
386  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
387  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
388 
389  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
390  ColorSpinorField &x, ColorSpinorField &b,
391  const QudaSolutionType) const;
392  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
393  const QudaSolutionType) const;
394  };
395 
396  // Full domain wall, but with 4-d parity ordered fields
397  class DiracDomainWall4D : public DiracDomainWall
398  {
399 
400 private:
401 public:
402  DiracDomainWall4D(const DiracParam &param);
403  DiracDomainWall4D(const DiracDomainWall4D &dirac);
404  virtual ~DiracDomainWall4D();
405  DiracDomainWall4D &operator=(const DiracDomainWall4D &dirac);
406 
407  void Dslash4(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
408  void Dslash5(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
409  void Dslash4Xpay(ColorSpinorField &out, const ColorSpinorField &in,
410  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
411  void Dslash5Xpay(ColorSpinorField &out, const ColorSpinorField &in,
412  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
413 
414  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
415  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
416 
417  void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b,
418  const QudaSolutionType) const;
419  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
420  };
421 
422  // 4d Even-odd preconditioned domain wall
423  class DiracDomainWall4DPC : public DiracDomainWall4D
424  {
425 
426 private:
427 public:
428  DiracDomainWall4DPC(const DiracParam &param);
429  DiracDomainWall4DPC(const DiracDomainWall4DPC &dirac);
430  virtual ~DiracDomainWall4DPC();
431  DiracDomainWall4DPC &operator=(const DiracDomainWall4DPC &dirac);
432 
433  void Dslash5inv(
434  ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const double &kappa5) const;
435  void Dslash5invXpay(ColorSpinorField &out, const ColorSpinorField &in,
436  const QudaParity parity, const double &kappa5, const ColorSpinorField &x, const double &k) const;
437 
438  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
439  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
440 
441  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
442  ColorSpinorField &x, ColorSpinorField &b,
443  const QudaSolutionType) const;
444  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
445  const QudaSolutionType) const;
446  };
447 
448  // Full Mobius
449  class DiracMobius : public DiracDomainWall {
450 
451  protected:
452  //Mobius coefficients
453  Complex b_5[QUDA_MAX_DWF_LS];
454  Complex c_5[QUDA_MAX_DWF_LS];
455 
461  bool zMobius;
462 
463  public:
464  DiracMobius(const DiracParam &param);
465  DiracMobius(const DiracMobius &dirac);
466  virtual ~DiracMobius();
467  DiracMobius& operator=(const DiracMobius &dirac);
468 
469  void Dslash4(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
470  void Dslash4pre(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
471  void Dslash5(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
472 
473  void Dslash4Xpay(ColorSpinorField &out, const ColorSpinorField &in,
474  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
475  void Dslash4preXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity,
476  const ColorSpinorField &x, const double &k) const;
477  void Dslash5Xpay(ColorSpinorField &out, const ColorSpinorField &in,
478  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
479 
480  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
481  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
482 
483  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
484  ColorSpinorField &x, ColorSpinorField &b,
485  const QudaSolutionType) const;
486  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
487  const QudaSolutionType) const;
488  };
489 
490  // 4d Even-odd preconditioned Mobius domain wall
491  class DiracMobiusPC : public DiracMobius {
492 
493  protected:
494 
495  private:
496 
497  public:
498  DiracMobiusPC(const DiracParam &param);
499  DiracMobiusPC(const DiracMobiusPC &dirac);
500  virtual ~DiracMobiusPC();
501  DiracMobiusPC& operator=(const DiracMobiusPC &dirac);
502 
503  void Dslash5inv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
504 
505  void Dslash5invXpay(ColorSpinorField &out, const ColorSpinorField &in,
506  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
507 
508 
509  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
510  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
511  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol, ColorSpinorField &x,
512  ColorSpinorField &b, const QudaSolutionType) const;
513  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
514  };
515 
516  // Full twisted mass
517  class DiracTwistedMass : public DiracWilson {
518 
519  protected:
520  mutable double mu;
521  mutable double epsilon;
522  void twistedApply(ColorSpinorField &out, const ColorSpinorField &in, const QudaTwistGamma5Type twistType) const;
523  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, QudaParity parity) const;
524  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, QudaParity parity,
525  const ColorSpinorField &x, const double &k) const;
526 
527  public:
528  DiracTwistedMass(const DiracTwistedMass &dirac);
529  DiracTwistedMass(const DiracParam &param, const int nDim);
530  virtual ~DiracTwistedMass();
531  DiracTwistedMass& operator=(const DiracTwistedMass &dirac);
532 
533  void Twist(ColorSpinorField &out, const ColorSpinorField &in) const;
534 
535  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
536  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
537 
538  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
539  ColorSpinorField &x, ColorSpinorField &b,
540  const QudaSolutionType) const;
541  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
542  const QudaSolutionType) const;
543 
544  double Mu() const { return mu; }
545 
557  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
558  double kappa, double mass, double mu, double mu_factor=0.) const;
559  };
560 
561  // Even-odd preconditioned twisted mass
562  class DiracTwistedMassPC : public DiracTwistedMass {
563 
564  public:
565  DiracTwistedMassPC(const DiracTwistedMassPC &dirac);
566  DiracTwistedMassPC(const DiracParam &param, const int nDim);
567 
568  virtual ~DiracTwistedMassPC();
569  DiracTwistedMassPC& operator=(const DiracTwistedMassPC &dirac);
570 
571  void TwistInv(ColorSpinorField &out, const ColorSpinorField &in) const;
572 
573  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
574  const QudaParity parity) const;
575  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
576  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
577  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
578  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
579 
580  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
581  ColorSpinorField &x, ColorSpinorField &b,
582  const QudaSolutionType) const;
583  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
584  const QudaSolutionType) const;
596  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
597  double kappa, double mass, double mu, double mu_factor=0.) const;
598  };
599 
600  // Full twisted mass with a clover term
601  class DiracTwistedClover : public DiracWilson {
602 
603  protected:
604  double mu;
605  double epsilon;
606  cudaCloverField &clover;
607  void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
608  void twistedCloverApply(ColorSpinorField &out, const ColorSpinorField &in,
609  const QudaTwistGamma5Type twistType, const int parity) const;
610 
611  public:
612  DiracTwistedClover(const DiracTwistedClover &dirac);
613  DiracTwistedClover(const DiracParam &param, const int nDim);
614  virtual ~DiracTwistedClover();
615  DiracTwistedClover& operator=(const DiracTwistedClover &dirac);
616 
617  void TwistClover(ColorSpinorField &out, const ColorSpinorField &in, const int parity) const;
618 
619  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
620  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity,
621  const ColorSpinorField &x, const double &k) const;
622 
623  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
624  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
625 
626  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
627  ColorSpinorField &x, ColorSpinorField &b,
628  const QudaSolutionType) const;
629  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
630  const QudaSolutionType) const;
631 
632  double Mu() const { return mu; }
633 
645  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
646  double kappa, double mass, double mu, double mu_factor=0.) const;
647  };
648 
649  // Even-odd preconditioned twisted mass with a clover term
650  class DiracTwistedCloverPC : public DiracTwistedClover {
651 
652  mutable bool reverse;
654 public:
655  DiracTwistedCloverPC(const DiracTwistedCloverPC &dirac);
656  DiracTwistedCloverPC(const DiracParam &param, const int nDim);
657 
658  virtual ~DiracTwistedCloverPC();
659  DiracTwistedCloverPC& operator=(const DiracTwistedCloverPC &dirac);
660 
661  void TwistCloverInv(ColorSpinorField &out, const ColorSpinorField &in, const int parity) const;
662 
663  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
664  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity,
665  const ColorSpinorField &x, const double &k) const;
666  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
667  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
668 
669  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
670  ColorSpinorField &x, ColorSpinorField &b,
671  const QudaSolutionType) const;
672  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
673  const QudaSolutionType) const;
674 
689  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
690  double kappa, double mass, double mu, double mu_factor=0.) const;
691  };
692 
693  // Full staggered
694  class DiracStaggered : public Dirac {
695 
696  protected:
697 
698  public:
699  DiracStaggered(const DiracParam &param);
700  DiracStaggered(const DiracStaggered &dirac);
701  virtual ~DiracStaggered();
702  DiracStaggered& operator=(const DiracStaggered &dirac);
703 
704  virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
705 
706  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
707  const QudaParity parity) const;
708  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
709  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
710  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
711  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
712 
713  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
714  ColorSpinorField &x, ColorSpinorField &b,
715  const QudaSolutionType) const;
716  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
717  const QudaSolutionType) const;
718 
730  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
731  double kappa, double mass, double mu=0., double mu_factor=0.) const;
732  };
733 
734  // Even-odd preconditioned staggered
735  class DiracStaggeredPC : public DiracStaggered {
736 
737  protected:
738 
739  public:
740  DiracStaggeredPC(const DiracParam &param);
741  DiracStaggeredPC(const DiracStaggeredPC &dirac);
742  virtual ~DiracStaggeredPC();
743  DiracStaggeredPC& operator=(const DiracStaggeredPC &dirac);
744 
745  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
746  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
747 
748  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
749  ColorSpinorField &x, ColorSpinorField &b,
750  const QudaSolutionType) const;
751  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
752  const QudaSolutionType) const;
753  };
754 
755  // Full staggered
756  class DiracImprovedStaggered : public Dirac {
757 
758  protected:
759  cudaGaugeField &fatGauge;
760  cudaGaugeField &longGauge;
761 
762  public:
763  DiracImprovedStaggered(const DiracParam &param);
764  DiracImprovedStaggered(const DiracImprovedStaggered &dirac);
765  virtual ~DiracImprovedStaggered();
766  DiracImprovedStaggered& operator=(const DiracImprovedStaggered &dirac);
767 
768  virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
769 
770  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
771  const QudaParity parity) const;
772  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
773  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
774  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
775  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
776 
777  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
778  ColorSpinorField &x, ColorSpinorField &b,
779  const QudaSolutionType) const;
780  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
781  const QudaSolutionType) const;
782  };
783 
784  // Even-odd preconditioned staggered
785  class DiracImprovedStaggeredPC : public DiracImprovedStaggered {
786 
787  protected:
788 
789  public:
790  DiracImprovedStaggeredPC(const DiracParam &param);
791  DiracImprovedStaggeredPC(const DiracImprovedStaggeredPC &dirac);
792  virtual ~DiracImprovedStaggeredPC();
793  DiracImprovedStaggeredPC& operator=(const DiracImprovedStaggeredPC &dirac);
794 
795  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
796  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
797 
798  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
799  ColorSpinorField &x, ColorSpinorField &b,
800  const QudaSolutionType) const;
801  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
802  const QudaSolutionType) const;
803  };
804 
809  class DiracCoarse : public Dirac {
810 
811  protected:
812  double mu;
813  double mu_factor;
814  const Transfer *transfer;
815  const Dirac *dirac;
816  const bool need_bidirectional;
818  mutable cpuGaugeField *Y_h;
819  mutable cpuGaugeField *X_h;
820  mutable cpuGaugeField *Xinv_h;
821  mutable cpuGaugeField *Yhat_h;
823  mutable cudaGaugeField *Y_d;
824  mutable cudaGaugeField *X_d;
825  mutable cudaGaugeField *Xinv_d;
826  mutable cudaGaugeField *Yhat_d;
832  void initializeCoarse();
833 
839  void initializeLazy(QudaFieldLocation location) const;
840 
841  mutable bool enable_gpu;
842  mutable bool enable_cpu;
843  const bool gpu_setup;
844  mutable bool init_gpu;
845  mutable bool init_cpu;
846  const bool mapped;
853  void createY(bool gpu = true, bool mapped = false) const;
854 
859  void createYhat(bool gpu = true) const;
860 
861  public:
862  double Mu() const { return mu; }
863  double MuFactor() const { return mu_factor; }
864 
870  DiracCoarse(const DiracParam &param, bool gpu_setup=true, bool mapped=false);
871 
883  DiracCoarse(const DiracParam &param,
884  cpuGaugeField *Y_h, cpuGaugeField *X_h, cpuGaugeField *Xinv_h, cpuGaugeField *Yhat_h,
885  cudaGaugeField *Y_d=0, cudaGaugeField *X_d=0, cudaGaugeField *Xinv_d=0, cudaGaugeField *Yhat_d=0);
886 
891  DiracCoarse(const DiracCoarse &dirac, const DiracParam &param);
892  virtual ~DiracCoarse();
893 
900  void Clover(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
901 
908  void CloverInv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
909 
916  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
917  const QudaParity parity) const;
918 
925  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity,
926  const ColorSpinorField &x, const double &k) const;
927 
933  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
934 
935  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
936 
937  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol, ColorSpinorField &x, ColorSpinorField &b,
938  const QudaSolutionType) const;
939 
940  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
941 
953  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
954  double kappa, double mass, double mu, double mu_factor=0.) const;
955 
956 
965  void createPreconditionedCoarseOp(GaugeField &Yhat, GaugeField &Xinv, const GaugeField &Y, const GaugeField &X);
966 
967  };
968 
972  class DiracCoarsePC : public DiracCoarse {
973 
974  public:
979  DiracCoarsePC(const DiracParam &param, bool gpu_setup=true);
980 
985  DiracCoarsePC(const DiracCoarse &dirac, const DiracParam &param);
986 
987  virtual ~DiracCoarsePC();
988 
989  void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
990  void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity,
991  const ColorSpinorField &x, const double &k) const;
992  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
993  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
994  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol, ColorSpinorField &x, ColorSpinorField &b,
995  const QudaSolutionType) const;
996  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
997 
1012  void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T,
1013  double kappa, double mass, double mu, double mu_factor=0.) const;
1014  };
1015 
1016 
1022  class GaugeLaplace : public Dirac {
1023 
1024  public:
1025  GaugeLaplace(const DiracParam &param);
1026  GaugeLaplace(const GaugeLaplace &laplace);
1027 
1028  virtual ~GaugeLaplace();
1029  GaugeLaplace& operator=(const GaugeLaplace &laplace);
1030 
1031  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
1032  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
1033  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
1034  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
1035  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
1036 
1037  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
1038  ColorSpinorField &x, ColorSpinorField &b,
1039  const QudaSolutionType) const;
1040  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
1041  const QudaSolutionType) const;
1042  };
1043 
1047  class GaugeLaplacePC : public GaugeLaplace {
1048 
1049  public:
1050  GaugeLaplacePC(const DiracParam &param);
1051  GaugeLaplacePC(const GaugeLaplacePC &laplace);
1052  virtual ~GaugeLaplacePC();
1053  GaugeLaplacePC& operator=(const GaugeLaplacePC &laplace);
1054 
1055  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
1056  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
1057 
1058  void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
1059  ColorSpinorField &x, ColorSpinorField &b,
1060  const QudaSolutionType) const;
1061  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
1062  };
1063 
1069  class GaugeCovDev : public Dirac {
1070 
1071  public:
1072  GaugeCovDev(const DiracParam &param);
1073  GaugeCovDev(const GaugeCovDev &covDev);
1074 
1075  virtual ~GaugeCovDev();
1076  GaugeCovDev& operator=(const GaugeCovDev &covDev);
1077 
1078  virtual void DslashCD(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const int mu) const;
1079  virtual void MCD(ColorSpinorField &out, const ColorSpinorField &in, const int mu) const;
1080  virtual void MdagMCD(ColorSpinorField &out, const ColorSpinorField &in, const int mu) const;
1081 
1082 
1083  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
1084  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
1085  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
1086  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
1087  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
1088 
1089  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
1090  ColorSpinorField &x, ColorSpinorField &b,
1091  const QudaSolutionType) const;
1092  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
1093  const QudaSolutionType) const;
1094 
1095 
1096  };
1097 
1098 
1099  // Functor base class for applying a given Dirac matrix (M, MdagM, etc.)
1100  class DiracMatrix {
1101 
1102  protected:
1103  const Dirac *dirac;
1104 
1105  public:
1106  DiracMatrix(const Dirac &d) : dirac(&d), shift(0.0) { }
1107  DiracMatrix(const Dirac *d) : dirac(d), shift(0.0) { }
1108  DiracMatrix(const DiracMatrix &mat) : dirac(mat.dirac), shift(mat.shift) { }
1109  DiracMatrix(const DiracMatrix *mat) : dirac(mat->dirac), shift(mat->shift) { }
1110  virtual ~DiracMatrix() { }
1111 
1112  virtual void operator()(ColorSpinorField &out, const ColorSpinorField &in) const = 0;
1113  virtual void operator()(ColorSpinorField &out, const ColorSpinorField &in,
1114  ColorSpinorField &tmp) const = 0;
1115  virtual void operator()(ColorSpinorField &out, const ColorSpinorField &in,
1116  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const = 0;
1117 
1118 
1119  unsigned long long flops() const { return dirac->Flops(); }
1120 
1121 
1122  QudaMatPCType getMatPCType() const { return dirac->getMatPCType(); }
1123 
1124  virtual int getStencilSteps() const = 0;
1125 
1126  std::string Type() const { return typeid(*dirac).name(); }
1127 
1128  bool isStaggered() const {
1129  return (Type() == typeid(DiracStaggeredPC).name() ||
1130  Type() == typeid(DiracStaggered).name() ||
1131  Type() == typeid(DiracImprovedStaggeredPC).name() ||
1132  Type() == typeid(DiracImprovedStaggered).name()) ? true : false;
1133  }
1134 
1135  const Dirac *Expose() const { return dirac; }
1136 
1138  double shift;
1139  };
1140 
1141  class DiracM : public DiracMatrix {
1142 
1143  public:
1144  DiracM(const Dirac &d) : DiracMatrix(d) { }
1145  DiracM(const Dirac *d) : DiracMatrix(d) { }
1146 
1147  void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
1148  {
1149  dirac->M(out, in);
1150  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1151  }
1152 
1153  void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
1154  {
1155  bool reset1 = false;
1156  if (!dirac->tmp1) { dirac->tmp1 = &tmp; reset1 = true; }
1157  dirac->M(out, in);
1158  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1159  if (reset1) { dirac->tmp1 = NULL; reset1 = false; }
1160  }
1161 
1162  void operator()(ColorSpinorField &out, const ColorSpinorField &in,
1163  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1164  {
1165  bool reset1 = false;
1166  bool reset2 = false;
1167  if (!dirac->tmp1) { dirac->tmp1 = &Tmp1; reset1 = true; }
1168  if (!dirac->tmp2) { dirac->tmp2 = &Tmp2; reset2 = true; }
1169  dirac->M(out, in);
1170  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1171  if (reset2) { dirac->tmp2 = NULL; reset2 = false; }
1172  if (reset1) { dirac->tmp1 = NULL; reset1 = false; }
1173  }
1174 
1175  int getStencilSteps() const
1176  {
1177  return dirac->getStencilSteps();
1178  }
1179  };
1180 
1181  class DiracMdagM : public DiracMatrix {
1182 
1183  public:
1184  DiracMdagM(const Dirac &d) : DiracMatrix(d) { }
1185  DiracMdagM(const Dirac *d) : DiracMatrix(d) { }
1186 
1187 
1188 
1189  void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
1190  {
1191  dirac->MdagM(out, in);
1192  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1193  }
1194 
1195  void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
1196  {
1197  dirac->tmp1 = &tmp;
1198  dirac->MdagM(out, in);
1199  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1200  dirac->tmp1 = NULL;
1201  }
1202 
1203  void operator()(ColorSpinorField &out, const ColorSpinorField &in,
1204  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1205  {
1206  dirac->tmp1 = &Tmp1;
1207  dirac->tmp2 = &Tmp2;
1208  dirac->MdagM(out, in);
1209  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1210  dirac->tmp2 = NULL;
1211  dirac->tmp1 = NULL;
1212  }
1213 
1214  int getStencilSteps() const
1215  {
1216  return 2*dirac->getStencilSteps(); // 2 for M and M dagger
1217  }
1218  };
1219 
1220 
1221  class DiracMMdag : public DiracMatrix {
1222 
1223  public:
1224  DiracMMdag(const Dirac &d) : DiracMatrix(d) { }
1225  DiracMMdag(const Dirac *d) : DiracMatrix(d) { }
1226 
1227  void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
1228  {
1229  dirac->MMdag(out, in);
1230  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1231  }
1232 
1233  void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
1234  {
1235  dirac->tmp1 = &tmp;
1236  dirac->MMdag(out, in);
1237  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1238  dirac->tmp1 = NULL;
1239  }
1240 
1241  void operator()(ColorSpinorField &out, const ColorSpinorField &in,
1242  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1243  {
1244  dirac->tmp1 = &Tmp1;
1245  dirac->tmp2 = &Tmp2;
1246  dirac->MMdag(out, in);
1247  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1248  dirac->tmp2 = NULL;
1249  dirac->tmp1 = NULL;
1250  }
1251 
1252  int getStencilSteps() const
1253  {
1254  return 2*dirac->getStencilSteps(); // 2 for M and M dagger
1255  }
1256  };
1257 
1258  class DiracMdag : public DiracMatrix {
1259 
1260  public:
1261  DiracMdag(const Dirac &d) : DiracMatrix(d) { }
1262  DiracMdag(const Dirac *d) : DiracMatrix(d) { }
1263 
1264  void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
1265  {
1266  dirac->Mdag(out, in);
1267  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1268  }
1269 
1270  void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
1271  {
1272  dirac->tmp1 = &tmp;
1273  dirac->Mdag(out, in);
1274  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1275  dirac->tmp1 = NULL;
1276  }
1277 
1278  void operator()(ColorSpinorField &out, const ColorSpinorField &in,
1279  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1280  {
1281  dirac->tmp1 = &Tmp1;
1282  dirac->tmp2 = &Tmp2;
1283  dirac->Mdag(out, in);
1284  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1285  dirac->tmp2 = NULL;
1286  dirac->tmp1 = NULL;
1287  }
1288 
1289  int getStencilSteps() const
1290  {
1291  return dirac->getStencilSteps();
1292  }
1293  };
1294 
1295  class DiracDagger : public DiracMatrix {
1296 
1297  protected:
1298  const DiracMatrix &mat;
1299 
1300  public:
1301  DiracDagger(const DiracMatrix &mat) : DiracMatrix(mat), mat(mat) { }
1302  DiracDagger(const DiracMatrix *mat) : DiracMatrix(mat), mat(*mat) { }
1303 
1304  void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
1305  {
1306  dirac->flipDagger();
1307  mat(out, in);
1308  dirac->flipDagger();
1309  }
1310 
1311  void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
1312  {
1313  dirac->flipDagger();
1314  mat(out, in, tmp);
1315  dirac->flipDagger();
1316  }
1317 
1318  void operator()(ColorSpinorField &out, const ColorSpinorField &in,
1319  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1320  {
1321  dirac->flipDagger();
1322  mat(out, in, Tmp1, Tmp2);
1323  dirac->flipDagger();
1324  }
1325 
1326  int getStencilSteps() const
1327  {
1328  return mat.getStencilSteps();
1329  }
1330  };
1331 
1332 } // namespace quda
1333 
1334 #endif // _DIRAC_QUDA_H
quda::DiracDagger::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1304
quda::DiracTwistedMass::Mu
double Mu() const
Definition: dirac_quda.h:544
quda::GaugeLaplacePC
Even-odd preconditioned Gauge Laplace operator.
Definition: dirac_quda.h:1047
quda::Dirac::type
QudaDiracType type
Definition: dirac_quda.h:124
quda::DiracMdagM::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1189
quda::Dirac::flops
unsigned long long flops
Definition: dirac_quda.h:121
quda::setDiracSloppyParam
void setDiracSloppyParam(DiracParam &diracParam, QudaInvertParam *inv_param, bool pc)
Definition: interface_quda.cpp:1661
quda::Dirac::Kappa
double Kappa() const
Definition: dirac_quda.h:173
quda::DiracMatrix::DiracMatrix
DiracMatrix(const Dirac &d)
Definition: dirac_quda.h:1106
quda::DiracMdag::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1270
QudaPrecision
enum QudaPrecision_s QudaPrecision
quda::DiracDagger::mat
const DiracMatrix & mat
Definition: dirac_quda.h:1298
quda::DiracParam::need_bidirectional
bool need_bidirectional
Definition: dirac_quda.h:51
quda::Dirac::flipDagger
void flipDagger() const
Definition: dirac_quda.h:183
quda::DiracMdagM::getStencilSteps
int getStencilSteps() const
Definition: dirac_quda.h:1214
quda::DiracMatrix::shift
double shift
Shift term added onto operator (M/M^dag M/M M^dag + shift)
Definition: dirac_quda.h:1138
quda::DiracMatrix::~DiracMatrix
virtual ~DiracMatrix()
Definition: dirac_quda.h:1110
QUDA_INVALID_PRECISION
Definition: enum_quda.h:63
quda::DiracCoarse::X_h
cpuGaugeField * X_h
Definition: dirac_quda.h:819
quda::Dirac::createCoarseOp
virtual void createCoarseOp(GaugeField &Y, GaugeField &X, const Transfer &T, double kappa, double mass=0., double mu=0., double mu_factor=0.) const
Create the coarse operator (virtual parent)
Definition: dirac_quda.h:196
errorQuda
#define errorQuda(...)
Definition: util_quda.h:121
quda::DiracM::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1147
color_spinor_field.h
quda::Dirac::gauge
cudaGaugeField * gauge
Definition: dirac_quda.h:115
quda::DiracMdag::DiracMdag
DiracMdag(const Dirac &d)
Definition: dirac_quda.h:1261
quda::DiracParam::transfer
Transfer * transfer
Definition: dirac_quda.h:49
quda::DiracMdagM::DiracMdagM
DiracMdagM(const Dirac *d)
Definition: dirac_quda.h:1185
quda::DiracMdag::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1264
tmp
cudaColorSpinorField * tmp
Definition: covdev_test.cpp:44
quda::DiracDagger::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1318
quda::DiracMdag::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1278
quda::Dirac::profile
TimeProfile profile
Definition: dirac_quda.h:132
quda::Dirac::getMatPCType
QudaMatPCType getMatPCType() const
Definition: dirac_quda.h:180
quda::Dirac::getStencilSteps
int getStencilSteps() const
Definition: dirac.cpp:251
quda::DiracCoarse::need_bidirectional
const bool need_bidirectional
Definition: dirac_quda.h:816
quda::Dirac::HaloPrecision
QudaPrecision HaloPrecision() const
Definition: dirac_quda.h:200
quda::DiracDagger::DiracDagger
DiracDagger(const DiracMatrix *mat)
Definition: dirac_quda.h:1302
quda::DiracParam::c_5
Complex c_5[QUDA_MAX_DWF_LS]
Definition: dirac_quda.h:28
quda::Dirac::setHaloPrecision
void setHaloPrecision(QudaPrecision halo_precision_) const
Definition: dirac_quda.h:201
quda::DiracParam::mu_factor
double mu_factor
Definition: dirac_quda.h:38
quda::Dirac::Flops
unsigned long long Flops() const
Definition: dirac_quda.h:177
QUDA_DAG_YES
Definition: enum_quda.h:219
quda::Dirac::Mu
virtual double Mu() const
Definition: dirac_quda.h:174
QUDA_DAG_INVALID
Definition: enum_quda.h:220
quda::Dirac::MdagM
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const =0
quda::DiracParam::gauge
cudaGaugeField * gauge
Definition: dirac_quda.h:31
quda::DiracImprovedStaggered::fatGauge
cudaGaugeField & fatGauge
Definition: dirac_quda.h:759
quda::DiracMdagM::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1203
param
QudaGaugeParam param
Definition: pack_test.cpp:17
quda::DiracCoarse::Y_d
cudaGaugeField * Y_d
Definition: dirac_quda.h:823
quda::setDiracParam
void setDiracParam(DiracParam &diracParam, QudaInvertParam *inv_param, bool pc)
Definition: interface_quda.cpp:1562
quda::DiracParam::b_5
Complex b_5[QUDA_MAX_DWF_LS]
Definition: dirac_quda.h:27
inv_param
QudaInvertParam inv_param
Definition: covdev_test.cpp:37
quda::covDev
__device__ __host__ void covDev(Arg &arg, int idx, int parity)
Definition: covDev.cuh:119
quda::DiracMdag::getStencilSteps
int getStencilSteps() const
Definition: dirac_quda.h:1289
quda::DiracMMdag::getStencilSteps
int getStencilSteps() const
Definition: dirac_quda.h:1252
quda::DiracTwistedClover::clover
cudaCloverField & clover
Definition: dirac_quda.h:606
quda::DiracCoarse::transfer
const Transfer * transfer
Definition: dirac_quda.h:814
quda::Dirac::MMdag
void MMdag(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac.cpp:97
quda::DiracClover::clover
cudaCloverField & clover
Definition: dirac_quda.h:269
quda::DiracMatrix::DiracMatrix
DiracMatrix(const Dirac *d)
Definition: dirac_quda.h:1107
quda::DiracCoarse::gpu_setup
const bool gpu_setup
Definition: dirac_quda.h:843
quda::DiracParam::commDim
int commDim[QUDA_MAX_DIM]
Definition: dirac_quda.h:44
quda::blas::axpy
void axpy(double a, ColorSpinorField &x, ColorSpinorField &y)
Definition: blas_quda.h:35
quda::DiracParam::dagger
QudaDagType dagger
Definition: dirac_quda.h:30
in
cpuColorSpinorField * in
Definition: staggered_invert_test.cpp:98
quda::Dirac::Dagger
void Dagger(QudaDagType dag) const
Definition: dirac_quda.h:182
QudaMatPCType
enum QudaMatPCType_s QudaMatPCType
quda::DiracM::getStencilSteps
int getStencilSteps() const
Definition: dirac_quda.h:1175
quda::Dirac::mass
double mass
Definition: dirac_quda.h:117
quda::DiracDagger::getStencilSteps
int getStencilSteps() const
Definition: dirac_quda.h:1326
quda::DiracMMdag::DiracMMdag
DiracMMdag(const Dirac &d)
Definition: dirac_quda.h:1224
QudaSolutionType
enum QudaSolutionType_s QudaSolutionType
quda::DiracParam::type
QudaDiracType type
Definition: dirac_quda.h:22
quda::Dirac::dagger
QudaDagType dagger
Definition: dirac_quda.h:120
quda::DiracParam::matpcType
QudaMatPCType matpcType
Definition: dirac_quda.h:29
X
int X[4]
Definition: covdev_test.cpp:70
quda::Complex
std::complex< double > Complex
Definition: quda_internal.h:46
quda::DiracCoarse::X_d
cudaGaugeField * X_d
Definition: dirac_quda.h:824
QudaDagType
enum QudaDagType_s QudaDagType
QudaParity
enum QudaParity_s QudaParity
quda::Dirac::MuFactor
virtual double MuFactor() const
Definition: dirac_quda.h:175
quda::DiracDagger::DiracDagger
DiracDagger(const DiracMatrix &mat)
Definition: dirac_quda.h:1301
quda::DiracCoarse::Mu
double Mu() const
Definition: dirac_quda.h:862
clover_field.h
quda::DiracM::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1162
quda::DiracParam::halo_precision
QudaPrecision halo_precision
Definition: dirac_quda.h:46
quda::Dirac::kappa
double kappa
Definition: dirac_quda.h:116
quda::Dirac::matpcType
QudaMatPCType matpcType
Definition: dirac_quda.h:119
quda::DiracM::DiracM
DiracM(const Dirac &d)
Definition: dirac_quda.h:1144
quda::Dirac::setMass
void setMass(double mass)
Definition: dirac_quda.h:169
quda::DiracMMdag::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1227
quda::Dirac::halo_precision
QudaPrecision halo_precision
Definition: dirac_quda.h:125
quda::Dirac::setCommDim
void setCommDim(const int commDim_[QUDA_MAX_DIM]) const
Enable / disable communications for the Dirac operator.
Definition: dirac_quda.h:145
quda::DiracMatrix::getMatPCType
QudaMatPCType getMatPCType() const
Definition: dirac_quda.h:1122
quda::Dirac::Mdag
void Mdag(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac.cpp:90
blas_quda.h
quda::Dirac::M
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const =0
quda::DiracM::DiracM
DiracM(const Dirac *d)
Definition: dirac_quda.h:1145
quda::DiracMatrix::dirac
const Dirac * dirac
Definition: dirac_quda.h:1103
QudaFieldLocation
enum QudaFieldLocation_s QudaFieldLocation
quda::DiracCoarse::Xinv_d
cudaGaugeField * Xinv_d
Definition: dirac_quda.h:825
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:99
QUDA_INVALID_DIRAC
Definition: enum_quda.h:320
quda::DiracDagger::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1311
quda::DiracCoarse::Yhat_d
cudaGaugeField * Yhat_d
Definition: dirac_quda.h:826
quda::DiracM::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1153
quda::DiracCoarse::Xinv_h
cpuGaugeField * Xinv_h
Definition: dirac_quda.h:820
quda::DiracMMdag::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1233
quda::DiracParam::clover
cudaCloverField * clover
Definition: dirac_quda.h:35
quda::DiracMatrix::flops
unsigned long long flops() const
Definition: dirac_quda.h:1119
QUDA_MAX_DWF_LS
#define QUDA_MAX_DWF_LS
Maximum length of the Ls dimension for domain-wall fermions.
Definition: quda_constants.h:49
quda::DiracCoarse::Y_h
cpuGaugeField * Y_h
Definition: dirac_quda.h:818
quda::DiracMdagM::DiracMdagM
DiracMdagM(const Dirac &d)
Definition: dirac_quda.h:1184
dslash_quda.h
quda::GaugeLaplace
Full Gauge Laplace operator. Although not a Dirac operator per se, it&#39;s a linear operator so it&#39;s con...
Definition: dirac_quda.h:1022
printfQuda
#define printfQuda(...)
Definition: util_quda.h:115
QUDA_DAG_NO
Definition: enum_quda.h:218
quda::DiracParam::fatGauge
cudaGaugeField * fatGauge
Definition: dirac_quda.h:32
quda::DiracMMdag::DiracMMdag
DiracMMdag(const Dirac *d)
Definition: dirac_quda.h:1225
quda::blas::flops
unsigned long long flops
Definition: blas_quda.cu:22
quda::DiracImprovedStaggered::longGauge
cudaGaugeField & longGauge
Definition: dirac_quda.h:760
quda::DiracCoarse::Yhat_h
cpuGaugeField * Yhat_h
Definition: dirac_quda.h:821
quda::DiracMatrix::getStencilSteps
virtual int getStencilSteps() const =0
quda::DiracParam::tmp2
ColorSpinorField * tmp2
Definition: dirac_quda.h:42
quda::DiracParam::longGauge
cudaGaugeField * longGauge
Definition: dirac_quda.h:33
quda::DiracMatrix::DiracMatrix
DiracMatrix(const DiracMatrix *mat)
Definition: dirac_quda.h:1109
quda::laplace
__device__ __host__ void laplace(Arg &arg, int idx, int parity)
Definition: laplace.cuh:132
QudaTwistGamma5Type
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
QUDA_MAX_DIM
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5...
Definition: quda_constants.h:17
quda::DiracCoarse::MuFactor
double MuFactor() const
Definition: dirac_quda.h:863
quda::DiracMatrix::DiracMatrix
DiracMatrix(const DiracMatrix &mat)
Definition: dirac_quda.h:1108
quda::DiracMdagM::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1195
QUDA_MATPC_INVALID
Definition: enum_quda.h:214
quda::DiracCoarse::dirac
const Dirac * dirac
Definition: dirac_quda.h:815
mat
void mat(void *out, void **link, void *in, int dagger_bit, int mu, QudaPrecision sPrecision, QudaPrecision gPrecision)
Definition: covdev_reference.cpp:112
quda::DiracMMdag::operator()
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1241
parity
QudaParity parity
Definition: covdev_test.cpp:54
quda::DiracMatrix::Type
std::string Type() const
Definition: dirac_quda.h:1126
QudaInvertParam_s
Definition: quda.h:97
quda::GaugeCovDev
Full Covariant Derivative operator. Although not a Dirac operator per se, it&#39;s a linear operator so i...
Definition: dirac_quda.h:1069
quda::DiracMdag::DiracMdag
DiracMdag(const Dirac *d)
Definition: dirac_quda.h:1262
quda::Dirac::tmp2
ColorSpinorField * tmp2
Definition: dirac_quda.h:123
gauge_field.h
quda::DiracMatrix::Expose
const Dirac * Expose() const
Definition: dirac_quda.h:1135
quda::DiracParam::tmp1
ColorSpinorField * tmp1
Definition: dirac_quda.h:41
quda::DiracMatrix::isStaggered
bool isStaggered() const
Definition: dirac_quda.h:1128
kappa5
double kappa5
Definition: dslash_ctest.cpp:31
quda::DiracCoarse::mapped
const bool mapped
Definition: dirac_quda.h:846
QudaDiracType
enum QudaDiracType_s QudaDiracType
quda_internal.h
quda::Dirac::tmp1
ColorSpinorField * tmp1
Definition: dirac_quda.h:122
Generated by   doxygen 1.8.13