8 #include <dslash_init.cuh>
11 namespace ndegtwisted {
12 #include <dslash_init.cuh>
15 namespace dslash_aux {
16 #include <dslash_init.cuh>
59 errorQuda(
"DiracTwistedMass::twistedApply method for flavor doublet is not implemented..\n");
93 twistedMassDslashCuda(&out.
Odd(),
gauge, &in.
Even(),
QUDA_ODD_PARITY,
dagger, &in.
Odd(),
QUDA_DEG_DSLASH_TWIST_XPAY, a, -
kappa, 0.0, 0.0,
commDim,
profile);
94 twistedMassDslashCuda(&out.
Even(),
gauge, &in.
Odd(),
QUDA_EVEN_PARITY,
dagger, &in.
Even(),
QUDA_DEG_DSLASH_TWIST_XPAY, a, -
kappa, 0.0, 0.0,
commDim,
profile);
99 ndegTwistedMassDslashCuda(&out.
Odd(),
gauge, &in.
Even(),
QUDA_ODD_PARITY,
dagger, &in.
Odd(),
QUDA_NONDEG_DSLASH, a, b, 1.0, -
kappa,
commDim,
profile);
100 ndegTwistedMassDslashCuda(&out.
Even(),
gauge, &in.
Odd(),
QUDA_EVEN_PARITY,
dagger, &in.
Even(),
QUDA_NONDEG_DSLASH, a, b, 1.0, -
kappa,
commDim,
profile);
123 errorQuda(
"Preconditioned solution requires a preconditioned solve_type");
148 if (&dirac !=
this) {
165 checkParitySpinor(in, out);
166 checkSpinorAlias(in, out);
173 twisted::setFace(face1,face2);
174 ndegtwisted::setFace(face1,face2);
178 double b = 1.0 / (1.0 + a*a);
180 twistedMassDslashCuda(&out,
gauge, &in, parity,
dagger, 0,
QUDA_DEG_DSLASH_TWIST_INV, a, b, 0.0, 0.0,
commDim, profile);
181 flops += 1392ll*in.
Volume();
183 twistedMassDslashCuda(&out,
gauge, &in, parity,
dagger, 0,
QUDA_DEG_TWIST_INV_DSLASH, a, b, 0.0, 0.0,
commDim, profile);
184 flops += 1392ll*in.
Volume();
188 double b = 2.0 *
kappa * epsilon;
189 double c = 1.0 / (1.0 + a*a - b*b);
192 ndegTwistedMassDslashCuda(&out,
gauge, &in, parity,
dagger, 0,
QUDA_NONDEG_DSLASH, a, b, c, 0.0,
commDim, profile);
193 flops += (1320ll+120ll)*in.
Volume();
196 bool reset = newTmp(&doubletTmp, in);
199 ndegTwistedMassDslashCuda(&out,
gauge, doubletTmp, parity,
dagger, 0,
QUDA_NONDEG_DSLASH, 0.0, 0.0, 1.0, 0.0,
commDim, profile);
201 flops += 1440ll*in.
Volume();
203 deleteTmp(&doubletTmp, reset);
212 checkParitySpinor(in, out);
213 checkSpinorAlias(in, out);
219 twisted::setFace(face1,face2);
220 ndegtwisted::setFace(face1,face2);
224 double b = k / (1.0 + a*a);
226 twistedMassDslashCuda(&out,
gauge, &in, parity,
dagger, &x,
QUDA_DEG_DSLASH_TWIST_INV, a, b, 0.0, 0.0,
commDim, profile);
227 flops += 1416ll*in.
Volume();
229 twistedMassDslashCuda(&out,
gauge, &in, parity,
dagger, &x,
QUDA_DEG_TWIST_INV_DSLASH, a, b, 0.0, 0.0,
commDim, profile);
230 flops += 1416ll*in.
Volume();
234 double b = 2.0 *
kappa * epsilon;
235 double c = 1.0 / (1.0 + a*a - b*b);
239 ndegTwistedMassDslashCuda(&out,
gauge, &in, parity,
dagger, &x,
QUDA_NONDEG_DSLASH, a, b, c, 0.0,
commDim, profile);
240 flops += 1464ll*in.
Volume();
243 bool reset = newTmp(&doubletTmp, in);
245 ndegTwistedMassDslashCuda(&out,
gauge, doubletTmp, parity,
dagger, &x,
QUDA_NONDEG_DSLASH, 0.0, 0.0, k, 0.0,
commDim, profile);
246 flops += 1464ll*in.
Volume();
247 deleteTmp(&doubletTmp, reset);
269 twistedMassDslashCuda(&out,
gauge,
tmp1,
QUDA_EVEN_PARITY,
dagger, &in,
QUDA_DEG_DSLASH_TWIST_XPAY, a, kappa2, 0.0, 0.0,
commDim,
profile);
273 twistedMassDslashCuda(&out,
gauge,
tmp1,
QUDA_ODD_PARITY,
dagger, &in,
QUDA_DEG_DSLASH_TWIST_XPAY, a, kappa2, 0.0, 0.0,
commDim,
profile);
289 double a = -2.0 * kappa *
mu;
290 double b = -2.0 * kappa *
epsilon;
295 ndegTwistedMassDslashCuda(&out,
gauge,
tmp1,
QUDA_EVEN_PARITY,
dagger, &in,
QUDA_NONDEG_DSLASH, a, b, c, kappa2,
commDim,
profile);
299 ndegTwistedMassDslashCuda(&out,
gauge,
tmp1,
QUDA_ODD_PARITY,
dagger, &in,
QUDA_NONDEG_DSLASH, a, b, c, kappa2,
commDim,
profile);
369 double d = (1.0 + a*a - bb*bb);
370 if(d <= 0)
errorQuda(
"Invalid twisted mass parameter\n");
379 ndegTwistedMassDslashCuda(
tmp1,
gauge, src,
QUDA_EVEN_PARITY,
dagger, &b.
Even(),
QUDA_NONDEG_DSLASH, 0.0, 0.0,
kappa, 0.0,
commDim,
profile);
389 ndegTwistedMassDslashCuda(
tmp1,
gauge, src,
QUDA_ODD_PARITY,
dagger, &b.
Odd(),
QUDA_NONDEG_DSLASH, 0.0, 0.0,
kappa, 0.0,
commDim,
profile);
398 ndegTwistedMassDslashCuda(src,
gauge,
tmp1,
QUDA_EVEN_PARITY,
dagger, &b.
Even(),
QUDA_NONDEG_DSLASH, 0.0, 0.0,
kappa, 0.0,
commDim,
profile);
407 ndegTwistedMassDslashCuda(src,
gauge,
tmp1,
QUDA_ODD_PARITY,
dagger, &b.
Odd(),
QUDA_NONDEG_DSLASH, 0.0, 0.0,
kappa, 0.0,
commDim,
profile);
450 double d = (1.0 + a*a - bb*bb);
451 if(d <= 0)
errorQuda(
"Invalid twisted mass parameter\n");
457 ndegTwistedMassDslashCuda(
tmp1,
gauge, &x.
Even(),
QUDA_ODD_PARITY,
dagger, &b.
Odd(),
QUDA_NONDEG_DSLASH, 0.0, 0.0,
kappa, 0.0,
commDim,
profile);
463 ndegTwistedMassDslashCuda(
tmp1,
gauge, &x.
Odd(),
QUDA_EVEN_PARITY,
dagger, &b.
Even(),
QUDA_NONDEG_DSLASH, 0.0, 0.0,
kappa, 0.0,
commDim,
profile);
virtual void checkParitySpinor(const cudaColorSpinorField &, const cudaColorSpinorField &) const
void prepare(cudaColorSpinorField *&src, cudaColorSpinorField *&sol, cudaColorSpinorField &x, cudaColorSpinorField &b, const QudaSolutionType) const
virtual ~DiracTwistedMassPC()
virtual void M(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
DiracTwistedMassPC(const DiracTwistedMassPC &dirac)
void TwistInv(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int mu
bool newTmp(cudaColorSpinorField **, const cudaColorSpinorField &) const
void reconstruct(cudaColorSpinorField &x, const cudaColorSpinorField &b, const QudaSolutionType) const
DiracTwistedMassPC & operator=(const DiracTwistedMassPC &dirac)
DiracWilson & operator=(const DiracWilson &dirac)
DiracTwistedMass & operator=(const DiracTwistedMass &dirac)
virtual void Dslash(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity) const
cudaColorSpinorField & Odd() const
int commDim[QUDA_MAX_DIM]
cudaColorSpinorField * tmp
void ndegTwistedMassDslashCuda(cudaColorSpinorField *out, const cudaGaugeField &gauge, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const QudaTwistDslashType type, const double &kappa, const double &mu, const double &epsilon, const double &k, const int *commDim, TimeProfile &profile, const QudaDslashPolicy &dslashPolicy=QUDA_DSLASH)
virtual void reconstruct(cudaColorSpinorField &x, const cudaColorSpinorField &b, const QudaSolutionType) const
VOLATILE spinorFloat kappa
enum QudaSolutionType_s QudaSolutionType
void deleteTmp(cudaColorSpinorField **, const bool &reset) const
enum QudaParity_s QudaParity
void twistGamma5Cuda(cudaColorSpinorField *out, const cudaColorSpinorField *in, const int dagger, const double &kappa, const double &mu, const double &epsilon, const QudaTwistGamma5Type)
ndeg tm:
virtual void prepare(cudaColorSpinorField *&src, cudaColorSpinorField *&sol, cudaColorSpinorField &x, cudaColorSpinorField &b, const QudaSolutionType) const
void MdagM(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
cudaColorSpinorField * tmp2
virtual void checkFullSpinor(const cudaColorSpinorField &, const cudaColorSpinorField &) const
virtual void DslashXpay(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity, const cudaColorSpinorField &x, const double &k) const
virtual void DslashXpay(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity, const cudaColorSpinorField &x, const double &k) const
cpuColorSpinorField * out
cudaColorSpinorField * tmp1
QudaTwistFlavorType TwistFlavor() const
void twistedApply(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaTwistGamma5Type twistType) const
virtual void MdagM(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
void M(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
DiracTwistedMass(const DiracTwistedMass &dirac)
void Mdag(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
void twistedMassDslashCuda(cudaColorSpinorField *out, const cudaGaugeField &gauge, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const QudaTwistDslashType type, const double &kappa, const double &mu, const double &epsilon, const double &k, const int *commDim, TimeProfile &profile, const QudaDslashPolicy &dslashPolicy=QUDA_DSLASH2)
QudaSiteSubset SiteSubset() const
virtual ~DiracTwistedMass()
cudaColorSpinorField & Even() const
void Twist(cudaColorSpinorField &out, const cudaColorSpinorField &in) const