QUDA  0.9.0
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:
23  double kappa;
24  double mass;
25  double m5; // used by domain wall only
26  int Ls;
27  double b_5[QUDA_MAX_DWF_LS];
28  double c_5[QUDA_MAX_DWF_LS];
32  cudaGaugeField *fatGauge; // used by staggered only
33  cudaGaugeField *longGauge; // used by staggered only
35 
36  double mu; // used by twisted mass only
37  double mu_factor; // used by multigrid only
38  double epsilon; //2nd tm parameter (used by twisted mass only)
39 
41  ColorSpinorField *tmp2; // used by Wilson-like kernels only
42 
43  int commDim[QUDA_MAX_DIM]; // whether to do comms or not
44 
45  // for multigrid only
48 
51  dagger(QUDA_DAG_INVALID), gauge(0), clover(0), mu(0.0), mu_factor(0.0), epsilon(0.0),
52  tmp1(0), tmp2(0)
53  {
54 
55  }
56 
57  void print() {
58  printfQuda("Printing DslashParam\n");
59  printfQuda("type = %d\n", type);
60  printfQuda("kappa = %g\n", kappa);
61  printfQuda("mass = %g\n", mass);
62  printfQuda("m5 = %g\n", m5);
63  printfQuda("Ls = %d\n", Ls);
64  printfQuda("matpcType = %d\n", matpcType);
65  printfQuda("dagger = %d\n", dagger);
66  printfQuda("mu = %g\n", mu);
67  printfQuda("epsilon = %g\n", epsilon);
68  for (int i=0; i<QUDA_MAX_DIM; i++) printfQuda("commDim[%d] = %d\n", i, commDim[i]);
69  for (int i=0; i<Ls; i++) printfQuda("b_5[%d] = %e\t c_5[%d] = %e\n", i,b_5[i],i,c_5[i]);
70  }
71  };
72 
73  void setDiracParam(DiracParam &diracParam, QudaInvertParam *inv_param, bool pc);
74  void setDiracSloppyParam(DiracParam &diracParam, QudaInvertParam *inv_param, bool pc);
75 
76  // forward declarations
77  class DiracMatrix;
78  class DiracM;
79  class DiracMdagM;
80  class DiracMMdag;
81  class DiracMdag;
82  //Forward declaration of multigrid Transfer class
83  class Transfer;
84 
85  // Abstract base class
86  class Dirac : public Object {
87 
88  friend class DiracMatrix;
89  friend class DiracM;
90  friend class DiracMdagM;
91  friend class DiracMMdag;
92  friend class DiracMdag;
93 
94  protected:
96  double kappa;
97  double mass;
99  mutable QudaDagType dagger; // mutable to simplify implementation of Mdag
100  mutable unsigned long long flops;
101  mutable ColorSpinorField *tmp1; // temporary hack
102  mutable ColorSpinorField *tmp2; // temporary hack
104 
105  bool newTmp(ColorSpinorField **, const ColorSpinorField &) const;
106  void deleteTmp(ColorSpinorField **, const bool &reset) const;
107 
109 
110  int commDim[QUDA_MAX_DIM]; // whether do comms or not
111 
113 
114  public:
115  Dirac(const DiracParam &param);
116  Dirac(const Dirac &dirac);
117  virtual ~Dirac();
118  Dirac& operator=(const Dirac &dirac);
119 
120  virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
121  virtual void checkFullSpinor(const ColorSpinorField &, const ColorSpinorField &) const;
122  void checkSpinorAlias(const ColorSpinorField &, const ColorSpinorField &) const;
123 
124  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
125  const QudaParity parity) const = 0;
126  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
127  const QudaParity parity, const ColorSpinorField &x,
128  const double &k) const = 0;
129  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const = 0;
130  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const = 0;
131  void Mdag(ColorSpinorField &out, const ColorSpinorField &in) const;
132  void MMdag(ColorSpinorField &out, const ColorSpinorField &in) const;
133 
134  // required methods to use e-o preconditioning for solving full system
135  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
137  const QudaSolutionType) const = 0;
138  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
139  const QudaSolutionType) const = 0;
140  void setMass(double mass){ this->mass = mass;}
141  // Dirac operator factory
142  static Dirac* create(const DiracParam &param);
143 
144  double Kappa() const { return kappa; }
145  virtual double Mu() const { return 0.; }
146  virtual double MuFactor() const { return 0.; }
147 
148  unsigned long long Flops() const { unsigned long long rtn = flops; flops = 0; return rtn; }
149 
150 
152  int getStencilSteps() const;
153  void Dagger(QudaDagType dag) const { dagger = dag; }
154 
167  virtual void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu=0., double mu_factor=0.) const
168  {errorQuda("Not implemented");}
169  };
170 
171  // Full Wilson
172  class DiracWilson : public Dirac {
173 
174  protected:
175  void initConstants();
176 
177  public:
178  DiracWilson(const DiracParam &param);
179  DiracWilson(const DiracWilson &dirac);
180  DiracWilson(const DiracParam &param, const int nDims);//to correctly adjust face for DW and non-deg twisted mass
181 
182  virtual ~DiracWilson();
184 
185  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
186  const QudaParity parity) const;
187  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
188  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
189  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
190  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
191 
192  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
194  const QudaSolutionType) const;
195  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
196  const QudaSolutionType) const;
197 
208  virtual void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu=0., double mu_factor=0.) const;
209  };
210 
211  // Even-odd preconditioned Wilson
212  class DiracWilsonPC : public DiracWilson {
213 
214  private:
215 
216  public:
219  virtual ~DiracWilsonPC();
221 
222  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
223  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
224 
227  const QudaSolutionType) const;
229  const QudaSolutionType) const;
230  };
231 
232  // Full clover
233  class DiracClover : public DiracWilson {
234 
235  protected:
237  void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
238  void initConstants();
239 
240  public:
241  DiracClover(const DiracParam &param);
242  DiracClover(const DiracClover &dirac);
243  virtual ~DiracClover();
245 
246  void Clover(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
247  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity,
248  const ColorSpinorField &x, const double &k) const;
249  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
250  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
251 
252  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
254  const QudaSolutionType) const;
255  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
256  const QudaSolutionType) const;
257 
268  void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu=0., double mu_factor=0.) const;
269  };
270 
271  // Even-odd preconditioned clover
272  class DiracCloverPC : public DiracClover {
273 
274  public:
277  virtual ~DiracCloverPC();
279 
280  void CloverInv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
282  const QudaParity parity) const;
284  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
285 
286  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
287  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
288 
291  const QudaSolutionType) const;
293  const QudaSolutionType) const;
294 
308  void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu=0., double mu_factor=0.) const;
309  };
310 
311 
312  // Full domain wall
313  class DiracDomainWall : public DiracWilson {
314 
315  protected:
316  double m5;
317  double kappa5;
318  int Ls; // length of the fifth dimension
319 
320  public:
323  virtual ~DiracDomainWall();
325 
327  const QudaParity parity) const;
329  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
330 
331  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
332  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
333 
334  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
336  const QudaSolutionType) const;
337  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
338  const QudaSolutionType) const;
339  };
340 
341  // 5d Even-odd preconditioned domain wall
343 
344  private:
345 
346  public:
349  virtual ~DiracDomainWallPC();
351 
352  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
353  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
354 
357  const QudaSolutionType) const;
359  const QudaSolutionType) const;
360  };
361 
362 // 4d Even-odd preconditioned domain wall
364 
365  private:
366 
367  public:
370  virtual ~DiracDomainWall4DPC();
372 
373  void Dslash4(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
374  void Dslash5(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
375  void Dslash5inv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const double &kappa5) const;
376 
378  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
380  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
382  const QudaParity parity, const double &kappa5, const ColorSpinorField &x, const double &k) const;
383 
384  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
385  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
386 
389  const QudaSolutionType) const;
391  const QudaSolutionType) const;
392  };
393 
394  // Full Mobius
395  class DiracMobius : public DiracDomainWall {
396 
397  protected:
398  //Mobius coefficients
401 
402  public:
403  DiracMobius(const DiracParam &param);
404  DiracMobius(const DiracMobius &dirac);
405  virtual ~DiracMobius();
407 
408  void Dslash4(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
409  void Dslash4pre(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
410  void Dslash5(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
411 
413  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
415  const ColorSpinorField &x, const double &k) const;
417  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
418 
419  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
420  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
421 
422  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
424  const QudaSolutionType) const;
425  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
426  const QudaSolutionType) const;
427  };
428 
429  // 4d Even-odd preconditioned Mobius domain wall
430  class DiracMobiusPC : public DiracMobius {
431 
432  protected:
433 
434  private:
435 
436  public:
439  virtual ~DiracMobiusPC();
441 
442  void Dslash5inv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
443 
445  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
446 
447 
448  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
449  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
451  ColorSpinorField &b, const QudaSolutionType) const;
452  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
453  };
454 
455  // Full twisted mass
456  class DiracTwistedMass : public DiracWilson {
457 
458  protected:
459  double mu;
460  double epsilon;
462  const QudaTwistGamma5Type twistType) const;
464  QudaParity parity, QudaTwistDslashType twistDslashType,
465  double a, double b, double c, double d) const;
468  QudaTwistDslashType twistDslashType,
469  double a, double b, double c, double d) const;
470 
472  QudaParity parity, QudaTwistDslashType twistDslashType,
473  double a, double b, double c, double d) const;
475  const ColorSpinorField &x, QudaParity parity, QudaTwistDslashType twistDslashType,
476  double a, double b, double c, double d) const;
477  public:
479  DiracTwistedMass(const DiracParam &param, const int nDim);
480  virtual ~DiracTwistedMass();
482 
483  void Twist(ColorSpinorField &out, const ColorSpinorField &in) const;
484 
485  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
486  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
487 
488  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
490  const QudaSolutionType) const;
491  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
492  const QudaSolutionType) const;
493 
494  double Mu() const { return mu; }
495 
508  void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const;
509  };
510 
511  // Even-odd preconditioned twisted mass
513 
514  public:
516  DiracTwistedMassPC(const DiracParam &param, const int nDim);
517 
518  virtual ~DiracTwistedMassPC();
520 
521  void TwistInv(ColorSpinorField &out, const ColorSpinorField &in) const;
522 
523  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
524  const QudaParity parity) const;
525  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
526  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
527  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
528  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
529 
532  const QudaSolutionType) const;
534  const QudaSolutionType) const;
547  void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const;
548  };
549 
550  // Full twisted mass with a clover term
552 
553  protected:
554  double mu;
555  double epsilon;
557  void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
559  const QudaTwistGamma5Type twistType, const int parity) const;
560 
561  public:
563  DiracTwistedClover(const DiracParam &param, const int nDim);
564  virtual ~DiracTwistedClover();
566 
567  void TwistClover(ColorSpinorField &out, const ColorSpinorField &in, const int parity) const; //IS PARITY REQUIRED???
568 
569  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
570  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
571 
572  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
574  const QudaSolutionType) const;
575  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
576  const QudaSolutionType) const;
577 
578  double Mu() const { return mu; }
579 
592  void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const;
593  };
594 
595  // Even-odd preconditioned twisted mass with a clover term
597 
598  public:
600  DiracTwistedCloverPC(const DiracParam &param, const int nDim);
601 
602  virtual ~DiracTwistedCloverPC();
604 
605  void TwistCloverInv(ColorSpinorField &out, const ColorSpinorField &in, const int parity) const;
606 
607  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
608  const QudaParity parity) const;
609  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
610  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
611  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
612  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
613 
616  const QudaSolutionType) const;
618  const QudaSolutionType) const;
619 
635  void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const;
636  };
637 
638  // Full staggered
639  class DiracStaggered : public Dirac {
640 
641  protected:
642 
643  public:
646  virtual ~DiracStaggered();
648 
649  virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
650 
651  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
652  const QudaParity parity) const;
653  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
654  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
655  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
656  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
657 
658  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
660  const QudaSolutionType) const;
661  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
662  const QudaSolutionType) const;
663  };
664 
665  // Even-odd preconditioned staggered
667 
668  protected:
669 
670  public:
673  virtual ~DiracStaggeredPC();
675 
676  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
677  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
678 
679  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
681  const QudaSolutionType) const;
682  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
683  const QudaSolutionType) const;
684  };
685 
686  // Full staggered
687  class DiracImprovedStaggered : public Dirac {
688 
689  protected:
692 
693  public:
696  virtual ~DiracImprovedStaggered();
698 
699  virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const;
700 
701  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
702  const QudaParity parity) const;
703  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
704  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
705  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
706  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
707 
708  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
710  const QudaSolutionType) const;
711  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
712  const QudaSolutionType) const;
713  };
714 
715  // Even-odd preconditioned staggered
717 
718  protected:
719 
720  public:
723  virtual ~DiracImprovedStaggeredPC();
725 
726  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
727  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
728 
729  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
731  const QudaSolutionType) const;
732  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
733  const QudaSolutionType) const;
734  };
735 
740  class DiracCoarse : public Dirac {
741 
742  protected:
743  double mu;
744  double mu_factor;
746  const Dirac *dirac;
758  void initializeCoarse();
760  bool enable_gpu;
761  bool init;
763  public:
764  double Mu() const { return mu; }
765  double MuFactor() const { return mu_factor; }
766 
771  DiracCoarse(const DiracParam &param, bool enable_gpu=true);
772 
787 
792  DiracCoarse(const DiracCoarse &dirac, const DiracParam &param);
793  virtual ~DiracCoarse();
794 
801  void Clover(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
802 
809  void CloverInv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
810 
817  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in,
818  const QudaParity parity) const;
819 
826  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity,
827  const ColorSpinorField &x, const double &k) const;
828 
834  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
835 
836  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
837 
839  const QudaSolutionType) const;
840 
841  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
842 
855  void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const;
856  };
857 
861  class DiracCoarsePC : public DiracCoarse {
862 
863  public:
864  DiracCoarsePC(const DiracParam &param, bool enable_gpu=true);
866  virtual ~DiracCoarsePC();
867 
868  void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
870  const ColorSpinorField &x, const double &k) const;
871  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
872  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
874  const QudaSolutionType) const;
875  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
876 
892  void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const;
893  };
894 
895 
901  class GaugeLaplace : public Dirac {
902 
903  public:
904  GaugeLaplace(const DiracParam &param);
906 
907  virtual ~GaugeLaplace();
909 
910  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
911  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
912  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
913  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
914  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
915 
916  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
918  const QudaSolutionType) const;
919  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
920  const QudaSolutionType) const;
921  };
922 
926  class GaugeLaplacePC : public GaugeLaplace {
927 
928  public:
931  virtual ~GaugeLaplacePC();
933 
934  void M(ColorSpinorField &out, const ColorSpinorField &in) const;
935  void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
936 
939  const QudaSolutionType) const;
940  void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const;
941  };
942 
948  class GaugeCovDev : public Dirac {
949 
950  public:
951  GaugeCovDev(const DiracParam &param);
953 
954  virtual ~GaugeCovDev();
956 
957  virtual void DslashCD(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const int mu) const;
958  virtual void MCD(ColorSpinorField &out, const ColorSpinorField &in, const int mu) const;
959  virtual void MdagMCD(ColorSpinorField &out, const ColorSpinorField &in, const int mu) const;
960 
961 
962  virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const;
963  virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in,
964  const QudaParity parity, const ColorSpinorField &x, const double &k) const;
965  virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const;
966  virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const;
967 
968  virtual void prepare(ColorSpinorField* &src, ColorSpinorField* &sol,
970  const QudaSolutionType) const;
971  virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b,
972  const QudaSolutionType) const;
973 
974 
975  };
976 
977 
978  // Functor base class for applying a given Dirac matrix (M, MdagM, etc.)
979  class DiracMatrix {
980 
981  protected:
982  const Dirac *dirac;
983 
984  public:
985  DiracMatrix(const Dirac &d) : dirac(&d) { }
986  DiracMatrix(const Dirac *d) : dirac(d) { }
987  virtual ~DiracMatrix() { }
988 
989  virtual void operator()(ColorSpinorField &out, const ColorSpinorField &in) const = 0;
990  virtual void operator()(ColorSpinorField &out, const ColorSpinorField &in,
991  ColorSpinorField &tmp) const = 0;
992  virtual void operator()(ColorSpinorField &out, const ColorSpinorField &in,
993  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const = 0;
994 
995  unsigned long long flops() const { return dirac->Flops(); }
996 
997 
999 
1000  virtual int getStencilSteps() const = 0;
1001 
1002  std::string Type() const { return typeid(*dirac).name(); }
1003 
1004  bool isStaggered() const {
1005  return (Type() == typeid(DiracStaggeredPC).name() ||
1006  Type() == typeid(DiracStaggered).name() ||
1007  Type() == typeid(DiracImprovedStaggeredPC).name() ||
1008  Type() == typeid(DiracImprovedStaggered).name()) ? true : false;
1009  }
1010 
1011  const Dirac* Expose() { return dirac; }
1012  };
1013 
1014  class DiracM : public DiracMatrix {
1015 
1016  public:
1017  DiracM(const Dirac &d) : DiracMatrix(d) { }
1018  DiracM(const Dirac *d) : DiracMatrix(d) { }
1019 
1021  {
1022  dirac->M(out, in);
1023  }
1024 
1026  {
1027  bool reset1 = false;
1028  if (!dirac->tmp1) { dirac->tmp1 = &tmp; reset1 = true; }
1029  dirac->M(out, in);
1030  if (reset1) { dirac->tmp1 = NULL; reset1 = false; }
1031  }
1032 
1034  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1035  {
1036  bool reset1 = false;
1037  bool reset2 = false;
1038  if (!dirac->tmp1) { dirac->tmp1 = &Tmp1; reset1 = true; }
1039  if (!dirac->tmp2) { dirac->tmp2 = &Tmp2; reset2 = true; }
1040  dirac->M(out, in);
1041  if (reset2) { dirac->tmp2 = NULL; reset2 = false; }
1042  if (reset1) { dirac->tmp1 = NULL; reset1 = false; }
1043  }
1044 
1045  int getStencilSteps() const
1046  {
1047  return dirac->getStencilSteps();
1048  }
1049  };
1050 
1051  class DiracMdagM : public DiracMatrix {
1052 
1053  public:
1054  DiracMdagM(const Dirac &d) : DiracMatrix(d), shift(0.0) { }
1055  DiracMdagM(const Dirac *d) : DiracMatrix(d), shift(0.0) { }
1056 
1058  double shift;
1059 
1061  {
1062  dirac->MdagM(out, in);
1063  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1064  }
1065 
1067  {
1068  dirac->tmp1 = &tmp;
1069  dirac->MdagM(out, in);
1070  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1071  dirac->tmp1 = NULL;
1072  }
1073 
1075  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1076  {
1077  dirac->tmp1 = &Tmp1;
1078  dirac->tmp2 = &Tmp2;
1079  dirac->MdagM(out, in);
1080  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1081  dirac->tmp2 = NULL;
1082  dirac->tmp1 = NULL;
1083  }
1084 
1085 
1086  int getStencilSteps() const
1087  {
1088  return 2*dirac->getStencilSteps(); // 2 for M and M dagger
1089  }
1090  };
1091 
1092 
1093  class DiracMMdag : public DiracMatrix {
1094 
1095  public:
1096  DiracMMdag(const Dirac &d) : DiracMatrix(d), shift(0.0) { }
1097  DiracMMdag(const Dirac *d) : DiracMatrix(d), shift(0.0) { }
1098 
1100  double shift;
1101 
1103  {
1104  dirac->MMdag(out, in);
1105  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1106  }
1107 
1109  {
1110  dirac->tmp1 = &tmp;
1111  dirac->MMdag(out, in);
1112  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1113  dirac->tmp1 = NULL;
1114  }
1115 
1117  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1118  {
1119  dirac->tmp1 = &Tmp1;
1120  dirac->tmp2 = &Tmp2;
1121  dirac->MMdag(out, in);
1122  if (shift != 0.0) blas::axpy(shift, const_cast<ColorSpinorField&>(in), out);
1123  dirac->tmp2 = NULL;
1124  dirac->tmp1 = NULL;
1125  }
1126 
1127 
1128  int getStencilSteps() const
1129  {
1130  return 2*dirac->getStencilSteps(); // 2 for M and M dagger
1131  }
1132  };
1133 
1134  class DiracMdag : public DiracMatrix {
1135 
1136  public:
1137  DiracMdag(const Dirac &d) : DiracMatrix(d) { }
1138  DiracMdag(const Dirac *d) : DiracMatrix(d) { }
1139 
1141  {
1142  dirac->Mdag(out, in);
1143  }
1144 
1146  {
1147  dirac->tmp1 = &tmp;
1148  dirac->Mdag(out, in);
1149  dirac->tmp1 = NULL;
1150  }
1151 
1153  ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
1154  {
1155  dirac->tmp1 = &Tmp1;
1156  dirac->tmp2 = &Tmp2;
1157  dirac->Mdag(out, in);
1158  dirac->tmp2 = NULL;
1159  dirac->tmp1 = NULL;
1160  }
1161 
1162  int getStencilSteps() const
1163  {
1164  return dirac->getStencilSteps();
1165  }
1166  };
1167 
1168 } // namespace quda
1169 
1170 #endif // _DIRAC_QUDA_H
DiracImprovedStaggeredPC & operator=(const DiracImprovedStaggeredPC &dirac)
virtual void MdagMCD(ColorSpinorField &out, const ColorSpinorField &in, const int mu) const
Dirac(const DiracParam &param)
Definition: dirac.cpp:12
double Mu() const
Definition: dirac_quda.h:494
Even-odd preconditioned Gauge Laplace operator.
Definition: dirac_quda.h:926
virtual ~DiracCoarse()
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
QudaDiracType type
Definition: dirac_quda.h:103
double c_5[QUDA_MAX_DWF_LS]
NEW: used by mobius domain wall only.
Definition: dirac_quda.h:28
void M(ColorSpinorField &out, const ColorSpinorField &in) const
DiracImprovedStaggeredPC(const DiracParam &param)
void M(ColorSpinorField &out, const ColorSpinorField &in) const
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1060
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
unsigned long long flops
Definition: dirac_quda.h:100
void setDiracSloppyParam(DiracParam &diracParam, QudaInvertParam *inv_param, bool pc)
DiracMobiusPC(const DiracParam &param)
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
double Kappa() const
Definition: dirac_quda.h:144
double mu
Definition: test_util.cpp:1643
DiracMatrix(const Dirac &d)
Definition: dirac_quda.h:985
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1145
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
const Dirac * Expose()
Definition: dirac_quda.h:1011
void Dslash5(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
void Dslash5(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
DiracWilson(const DiracParam &param)
Definition: dirac_wilson.cpp:8
int getStencilSteps() const
Definition: dirac_quda.h:1086
DiracDomainWall4DPC & operator=(const DiracDomainWall4DPC &dirac)
DiracWilsonPC & operator=(const DiracWilsonPC &dirac)
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
GaugeLaplacePC & operator=(const GaugeLaplacePC &laplace)
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
void TwistedDslash(ColorSpinorField &out, const ColorSpinorField &in, QudaParity parity, QudaTwistDslashType twistDslashType, double a, double b, double c, double d) const
virtual ~DiracMatrix()
Definition: dirac_quda.h:987
GaugeLaplacePC(const DiracParam &param)
cpuGaugeField * X_h
Definition: dirac_quda.h:749
const void * src
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
DiracTwistedMassPC(const DiracTwistedMassPC &dirac)
DiracWilsonPC(const DiracParam &param)
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
#define errorQuda(...)
Definition: util_quda.h:90
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1020
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
cudaGaugeField * gauge
Definition: dirac_quda.h:95
DiracStaggeredPC(const DiracParam &param)
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
void CloverInv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
Apply the inverse coarse clover operator.
DiracMdag(const Dirac &d)
Definition: dirac_quda.h:1137
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
double mu_factor[QUDA_MAX_MG_LEVEL]
Definition: test_util.cpp:1659
virtual void checkFullSpinor(const ColorSpinorField &, const ColorSpinorField &) const
Definition: dirac.cpp:129
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
void M(ColorSpinorField &out, const ColorSpinorField &in) const
Apply the full operator.
void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
double b_5[QUDA_MAX_DWF_LS]
Definition: dirac_quda.h:399
Transfer * transfer
Definition: dirac_quda.h:46
void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
Apply DslashXpay out = (D * in)
void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const
Create the coarse twisted-clover operator.
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
DiracMdagM(const Dirac *d)
Definition: dirac_quda.h:1055
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1140
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
GaugeLaplace(const DiracParam &param)
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
DiracTwistedMassPC & operator=(const DiracTwistedMassPC &dirac)
cudaColorSpinorField * tmp
Definition: covdev_test.cpp:44
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
void deleteTmp(ColorSpinorField **, const bool &reset) const
Definition: dirac.cpp:64
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1152
TimeProfile profile
Definition: dirac_quda.h:112
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
QudaMatPCType getMatPCType() const
Definition: dirac_quda.h:151
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
Apply DslashXpay out = (D * in + A * x)
int getStencilSteps() const
Definition: dirac.cpp:228
DiracWilson & operator=(const DiracWilson &dirac)
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
GaugeCovDev(const DiracParam &param)
Definition: gauge_covdev.cpp:9
void Clover(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
Apply the coarse clover operator.
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
virtual ~Dirac()
Definition: dirac.cpp:28
DiracTwistedMass & operator=(const DiracTwistedMass &dirac)
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const
Create the coarse even-odd preconditioned twisted-clover operator. Unlike the Wilson operator...
void M(ColorSpinorField &out, const ColorSpinorField &in) const
double mu_factor
Definition: dirac_quda.h:37
unsigned long long Flops() const
Definition: dirac_quda.h:148
void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const
Create the coarse even-odd preconditioned twisted-mass operator.
virtual double Mu() const
Definition: dirac_quda.h:145
void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const
Create the coarse operator from this coarse operator.
DiracTwistedCloverPC(const DiracTwistedCloverPC &dirac)
void Dslash5invXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const =0
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
virtual ~DiracClover()
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
cudaGaugeField * gauge
Definition: dirac_quda.h:31
__device__ __host__ void laplace(Arg &arg, int x_cb, int parity)
Definition: laplace.cu:113
void Twist(ColorSpinorField &out, const ColorSpinorField &in) const
cudaGaugeField & fatGauge
Definition: dirac_quda.h:690
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1074
void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const
void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
Apply DslashXpay out = (D * in + A * x)
QudaGaugeParam param
Definition: pack_test.cpp:17
cudaGaugeField * Y_d
Definition: dirac_quda.h:753
bool newTmp(ColorSpinorField **, const ColorSpinorField &) const
Definition: dirac.cpp:53
void M(ColorSpinorField &out, const ColorSpinorField &in) const
#define b
void setDiracParam(DiracParam &diracParam, QudaInvertParam *inv_param, bool pc)
DiracStaggered & operator=(const DiracStaggered &dirac)
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
void TwistCloverInv(ColorSpinorField &out, const ColorSpinorField &in, const int parity) const
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
int commDim[QUDA_MAX_DIM]
Definition: dirac_quda.h:110
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
Apply DslashXpay out = (D * in)
double shift
Shift term added onto operator (M^dag M + shift)
Definition: dirac_quda.h:1100
void Dslash4(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
void M(ColorSpinorField &out, const ColorSpinorField &in) const
DiracDomainWallPC(const DiracParam &param)
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
DiracClover & operator=(const DiracClover &dirac)
QudaInvertParam inv_param
Definition: covdev_test.cpp:37
DiracMobius(const DiracParam &param)
Definition: dirac_mobius.cpp:7
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const
Create the coarse even-odd preconditioned coarse operator. Unlike the Wilson operator, the coarsening of the preconditioned coarse operator differs from that of the unpreconditioned coarse operator, so we need to specialize it.
DiracStaggeredPC & operator=(const DiracStaggeredPC &dirac)
int getStencilSteps() const
Definition: dirac_quda.h:1162
int getStencilSteps() const
Definition: dirac_quda.h:1128
cudaCloverField & clover
Definition: dirac_quda.h:556
const Transfer * transfer
Definition: dirac_quda.h:745
void MMdag(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac.cpp:80
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
void Dslash5inv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const double &kappa5) const
void Dslash5invXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const double &kappa5, const ColorSpinorField &x, const double &k) const
DiracDomainWall & operator=(const DiracDomainWall &dirac)
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
void Clover(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
virtual void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu=0., double mu_factor=0.) const
Create the coarse operator (virtual parent)
Definition: dirac_quda.h:167
cudaCloverField & clover
Definition: dirac_quda.h:236
DiracMatrix(const Dirac *d)
Definition: dirac_quda.h:986
DiracTwistedCloverPC & operator=(const DiracTwistedCloverPC &dirac)
void checkSpinorAlias(const ColorSpinorField &, const ColorSpinorField &) const
Definition: dirac.cpp:137
DiracDomainWall4DPC(const DiracParam &param)
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const =0
int commDim[QUDA_MAX_DIM]
Definition: dirac_quda.h:43
QudaDagType dagger
Definition: dirac_quda.h:30
cpuColorSpinorField * in
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const =0
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
void Dagger(QudaDagType dag) const
Definition: dirac_quda.h:153
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
void TwistedDslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const ColorSpinorField &x, QudaParity parity, QudaTwistDslashType twistDslashType, double a, double b, double c, double d) const
enum QudaMatPCType_s QudaMatPCType
int getStencilSteps() const
Definition: dirac_quda.h:1045
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
DiracTwistedClover & operator=(const DiracTwistedClover &dirac)
double mass
Definition: dirac_quda.h:97
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
QudaTune tune
Definition: dirac_quda.h:108
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
DiracDomainWall(const DiracParam &param)
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
DiracMMdag(const Dirac &d)
Definition: dirac_quda.h:1096
enum QudaSolutionType_s QudaSolutionType
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
double b_5[QUDA_MAX_DWF_LS]
NEW: used by domain wall and twisted mass.
Definition: dirac_quda.h:27
QudaDiracType type
Definition: dirac_quda.h:22
DiracTwistedClover(const DiracTwistedClover &dirac)
virtual void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu=0., double mu_factor=0.) const
Create the coarse Wilson operator.
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
void Dslash4Xpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
void Dslash4pre(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const
QudaDagType dagger
Definition: dirac_quda.h:99
QudaMatPCType matpcType
NEW: used by mobius domain wall only.
Definition: dirac_quda.h:29
cudaGaugeField * X_d
Definition: dirac_quda.h:754
enum QudaDagType_s QudaDagType
void Dslash4preXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
enum QudaParity_s QudaParity
GaugeCovDev & operator=(const GaugeCovDev &covDev)
virtual double MuFactor() const
Definition: dirac_quda.h:146
void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu=0., double mu_factor=0.) const
Create the coarse even-odd preconditioned clover operator. Unlike the Wilson operator, the coarsening of the preconditioned clover operator differs from that of the unpreconditioned clover operator, so we need to specialize it.
void TwistClover(ColorSpinorField &out, const ColorSpinorField &in, const int parity) const
void Dslash4Xpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
DiracCloverPC(const DiracParam &param)
double Mu() const
Definition: dirac_quda.h:764
void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu=0., double mu_factor=0.) const
Create the coarse clover operator.
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1033
void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
DiracDomainWallPC & operator=(const DiracDomainWallPC &dirac)
double kappa
Definition: dirac_quda.h:96
QudaMatPCType matpcType
Definition: dirac_quda.h:98
double shift
Shift term added onto operator (M^dag M + shift)
Definition: dirac_quda.h:1058
DiracM(const Dirac &d)
Definition: dirac_quda.h:1017
void setMass(double mass)
Definition: dirac_quda.h:140
void operator()(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac_quda.h:1102
void Dslash4(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
void NdegTwistedDslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const ColorSpinorField &x, QudaParity parity, QudaTwistDslashType twistDslashType, double a, double b, double c, double d) const
void axpy(const double &a, ColorSpinorField &x, ColorSpinorField &y)
Definition: blas_quda.cu:150
QudaMatPCType getMatPCType() const
Definition: dirac_quda.h:998
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
enum QudaTwistDslashType_s QudaTwistDslashType
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
DiracClover(const DiracParam &param)
Definition: dirac_clover.cpp:8
void Mdag(ColorSpinorField &out, const ColorSpinorField &in) const
Definition: dirac.cpp:73
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const =0
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
DiracM(const Dirac *d)
Definition: dirac_quda.h:1018
void Dslash5Xpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
const Dirac * dirac
Definition: dirac_quda.h:982
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
Dirac & operator=(const Dirac &dirac)
Definition: dirac.cpp:32
double c_5[QUDA_MAX_DWF_LS]
Definition: dirac_quda.h:400
cudaGaugeField * Xinv_d
Definition: dirac_quda.h:755
GaugeCovDev * dirac
Definition: covdev_test.cpp:75
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
cpuColorSpinorField * out
void M(ColorSpinorField &out, const ColorSpinorField &in) const
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
cudaGaugeField * Yhat_d
Definition: dirac_quda.h:756
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1025
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
cpuGaugeField * Xinv_h
Definition: dirac_quda.h:750
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1108
cudaCloverField * clover
Definition: dirac_quda.h:34
unsigned long long flops() const
Definition: dirac_quda.h:995
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
#define QUDA_MAX_DWF_LS
Maximum length of the Ls dimension for domain-wall fermions.
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
cpuGaugeField * Y_h
Definition: dirac_quda.h:748
DiracMdagM(const Dirac &d)
Definition: dirac_quda.h:1054
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:901
virtual void operator()(ColorSpinorField &out, const ColorSpinorField &in) const =0
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
#define printfQuda(...)
Definition: util_quda.h:84
enum QudaTune_s QudaTune
cudaGaugeField * fatGauge
Definition: dirac_quda.h:32
virtual ~GaugeCovDev()
void M(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
DiracMMdag(const Dirac *d)
Definition: dirac_quda.h:1097
virtual void DslashCD(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const int mu) const
void NdegTwistedDslash(ColorSpinorField &out, const ColorSpinorField &in, QudaParity parity, QudaTwistDslashType twistDslashType, double a, double b, double c, double d) const
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
cudaGaugeField & longGauge
Definition: dirac_quda.h:691
void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
cpuGaugeField * Yhat_h
Definition: dirac_quda.h:751
virtual int getStencilSteps() const =0
void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
const void * c
ColorSpinorField * tmp2
Definition: dirac_quda.h:41
void createCoarseOp(GaugeField &Y, GaugeField &X, GaugeField &Xinv, GaugeField &Yhat, const Transfer &T, double kappa, double mu, double mu_factor=0.) const
Create the coarse twisted-mass operator.
void Dslash5Xpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
cudaGaugeField * longGauge
Definition: dirac_quda.h:33
virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
GaugeLaplace & operator=(const GaugeLaplace &laplace)
void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5...
double MuFactor() const
Definition: dirac_quda.h:765
void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &tmp) const
Definition: dirac_quda.h:1066
DiracTwistedMass(const DiracTwistedMass &dirac)
static Dirac * create(const DiracParam &param)
Definition: dirac.cpp:142
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual ~DiracMobius()
void TwistInv(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
Apply the full operator.
virtual void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const
Definition: dirac.cpp:89
const Dirac * dirac
Definition: dirac_quda.h:746
virtual ~DiracWilson()
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
DiracStaggered(const DiracParam &param)
DiracCloverPC & operator=(const DiracCloverPC &dirac)
void Dslash5inv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
void twistedCloverApply(ColorSpinorField &out, const ColorSpinorField &in, const QudaTwistGamma5Type twistType, const int parity) const
void CloverInv(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const =0
void operator()(ColorSpinorField &out, const ColorSpinorField &in, ColorSpinorField &Tmp1, ColorSpinorField &Tmp2) const
Definition: dirac_quda.h:1116
static __inline__ size_t size_t d
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const =0
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
DiracMobiusPC & operator=(const DiracMobiusPC &dirac)
QudaParity parity
Definition: covdev_test.cpp:53
std::string Type() const
Definition: dirac_quda.h:1002
virtual void Dslash(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity) const
Full Covariant Derivative operator. Although not a Dirac operator per se, it&#39;s a linear operator so i...
Definition: dirac_quda.h:948
DiracCoarse(const DiracParam &param, bool enable_gpu=true)
Definition: dirac_coarse.cpp:7
DiracMdag(const Dirac *d)
Definition: dirac_quda.h:1138
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const
ColorSpinorField * tmp2
Definition: dirac_quda.h:102
void M(ColorSpinorField &out, const ColorSpinorField &in) const
#define a
DiracImprovedStaggered(const DiracParam &param)
ColorSpinorField * tmp1
Definition: dirac_quda.h:40
bool isStaggered() const
Definition: dirac_quda.h:1004
DiracMobius & operator=(const DiracMobius &dirac)
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const
DiracCoarsePC(const DiracParam &param, bool enable_gpu=true)
virtual void MCD(ColorSpinorField &out, const ColorSpinorField &in, const int mu) const
enum QudaDiracType_s QudaDiracType
void checkParitySpinor(const ColorSpinorField &, const ColorSpinorField &) const
void covDev(cudaColorSpinorField *out, cudaGaugeField &gauge, const cudaColorSpinorField *in, const int parity, const int mu, TimeProfile &profile)
void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const
void twistedApply(ColorSpinorField &out, const ColorSpinorField &in, const QudaTwistGamma5Type twistType) const
DiracImprovedStaggered & operator=(const DiracImprovedStaggered &dirac)
ColorSpinorField * tmp1
Definition: dirac_quda.h:101
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const