7 namespace twistedclover {
8 #include <dslash_init.cuh>
11 namespace dslash_aux {
12 #include <dslash_init.cuh>
16 :
DiracWilson(param, nDim),
mu(param.
mu), epsilon(param.epsilon), clover(*(param.clover)), cloverInv(*(param.cloverInv))
23 :
DiracWilson(dirac),
mu(dirac.
mu), epsilon(dirac.epsilon), clover(dirac.clover), cloverInv(dirac.cloverInv)
76 errorQuda(
"DiracTwistedClover::twistedCloverApply method for flavor doublet is not implemented..\n");
111 twistedCloverDslashCuda(&out.
Odd(),
gauge, cs, cI, &in.
Even(),
QUDA_ODD_PARITY,
dagger, &in.
Odd(),
QUDA_DEG_DSLASH_CLOVER_TWIST_XPAY, a, -
kappa, 0.0, 0.0,
commDim,
profile);
112 twistedCloverDslashCuda(&out.
Even(),
gauge, cs, cI, &in.
Odd(),
QUDA_EVEN_PARITY,
dagger, &in.
Even(),
QUDA_DEG_DSLASH_CLOVER_TWIST_XPAY, a, -
kappa, 0.0, 0.0,
commDim,
profile);
115 errorQuda(
"Non-deg twisted clover not implemented yet");
138 errorQuda(
"Preconditioned solution requires a preconditioned solve_type");
163 if (&dirac !=
this) {
180 checkParitySpinor(in, out);
181 checkSpinorAlias(in, out);
188 twistedclover::setFace(face1,face2);
197 twistedCloverDslashCuda(&out,
gauge, cs, cI, &in, parity,
dagger, 0,
QUDA_DEG_DSLASH_CLOVER_TWIST_INV, a, b, 0.0, 0.0,
commDim, profile);
198 flops += 2376ll*in.
Volume();
200 twistedCloverDslashCuda(&out,
gauge, cs, cI, &in, parity,
dagger, 0,
QUDA_DEG_CLOVER_TWIST_INV_DSLASH, a, b, 0.0, 0.0,
commDim, profile);
201 flops += 1320ll*in.
Volume();
204 errorQuda(
"Non-degenerate DiracTwistedCloverPC is not implemented \n");
214 checkParitySpinor(in, out);
215 checkSpinorAlias(in, out);
221 twistedclover::setFace(face1,face2);
231 twistedCloverDslashCuda(&out,
gauge, cs, cI, &in, parity,
dagger, &x,
QUDA_DEG_DSLASH_CLOVER_TWIST_INV, a, b, 0.0, 0.0,
commDim, profile);
232 flops += 2400ll*in.
Volume();
234 twistedCloverDslashCuda(&out,
gauge, cs, cI, &in, parity,
dagger, &x,
QUDA_DEG_CLOVER_TWIST_INV_DSLASH, a, b, 0.0, 0.0,
commDim, profile);
235 flops += 1344ll*in.
Volume();
238 errorQuda(
"Non-degenerate DiracTwistedCloverPC is not implemented \n");
278 twistedCloverDslashCuda(&out,
gauge, cs, cI,
tmp1,
QUDA_EVEN_PARITY,
dagger, &in,
QUDA_DEG_DSLASH_CLOVER_TWIST_XPAY, a, kappa2, 0.0, 0.0,
commDim,
profile);
282 twistedCloverDslashCuda(&out,
gauge, cs, cI,
tmp1,
QUDA_ODD_PARITY,
dagger, &in,
QUDA_DEG_DSLASH_CLOVER_TWIST_XPAY, a, kappa2, 0.0, 0.0,
commDim,
profile);
289 errorQuda(
"Non-degenerate DiracTwistedCloverPC is not implemented \n");
352 errorQuda(
"Non-degenrate DiracTwistedCloverPC is not implemented \n");
384 errorQuda(
"Non-degenrate DiracTwistedCloverPC is not implemented \n");
virtual void reconstruct(cudaColorSpinorField &x, const cudaColorSpinorField &b, const QudaSolutionType) const
virtual void checkParitySpinor(const cudaColorSpinorField &, const cudaColorSpinorField &) const
void reconstruct(cudaColorSpinorField &x, const cudaColorSpinorField &b, const QudaSolutionType) const
void twistedCloverApply(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaTwistGamma5Type twistType, const int parity) const
__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int mu
bool newTmp(cudaColorSpinorField **, const cudaColorSpinorField &) const
DiracWilson & operator=(const DiracWilson &dirac)
virtual void M(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
DiracTwistedCloverPC(const DiracTwistedCloverPC &dirac)
cudaColorSpinorField & Odd() const
int commDim[QUDA_MAX_DIM]
void TwistClover(cudaColorSpinorField &out, const cudaColorSpinorField &in, const int parity) const
cudaColorSpinorField * tmp
void M(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
VOLATILE spinorFloat kappa
DiracTwistedCloverPC & operator=(const DiracTwistedCloverPC &dirac)
virtual void MdagM(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
DiracTwistedClover & operator=(const DiracTwistedClover &dirac)
enum QudaSolutionType_s QudaSolutionType
virtual ~DiracTwistedCloverPC()
void twistedCloverDslashCuda(cudaColorSpinorField *out, const cudaGaugeField &gauge, const FullClover *clover, const FullClover *cloverInv, const cudaColorSpinorField *in, const int parity, const int dagger, const cudaColorSpinorField *x, const QudaTwistCloverDslashType type, const double &kappa, const double &mu, const double &epsilon, const double &k, const int *commDim, TimeProfile &profile, const QudaDslashPolicy &dslashPolicy=QUDA_DSLASH2)
void deleteTmp(cudaColorSpinorField **, const bool &reset) const
DiracTwistedClover(const DiracTwistedClover &dirac)
void prepare(cudaColorSpinorField *&src, cudaColorSpinorField *&sol, cudaColorSpinorField &x, cudaColorSpinorField &b, const QudaSolutionType) const
void MdagM(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
enum QudaParity_s QudaParity
virtual void DslashXpay(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity, const cudaColorSpinorField &x, const double &k) const
virtual void prepare(cudaColorSpinorField *&src, cudaColorSpinorField *&sol, cudaColorSpinorField &x, cudaColorSpinorField &b, const QudaSolutionType) const
virtual ~DiracTwistedClover()
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
void twistCloverGamma5Cuda(cudaColorSpinorField *out, const cudaColorSpinorField *in, const int dagger, const double &kappa, const double &mu, const double &epsilon, const QudaTwistGamma5Type twist, const FullClover *clov, const FullClover *clovInv, const int parity)
cpuColorSpinorField * out
void TwistCloverInv(cudaColorSpinorField &out, const cudaColorSpinorField &in, const int parity) const
cudaColorSpinorField * tmp1
QudaTwistFlavorType TwistFlavor() const
cudaCloverField & cloverInv
enum QudaTwistGamma5Type_s QudaTwistGamma5Type
virtual void Dslash(cudaColorSpinorField &out, const cudaColorSpinorField &in, const QudaParity parity) const
void Mdag(cudaColorSpinorField &out, const cudaColorSpinorField &in) const
QudaSiteSubset SiteSubset() const
cudaColorSpinorField & Even() const
void checkParitySpinor(const cudaColorSpinorField &, const cudaColorSpinorField &) const