21 template <
typename Float,
int nDim,
int nColor,
int nParity,
bool dagger,
bool xpay, KernelType kernel_type,
typename Arg>
23 static constexpr
const char *
kernel =
"quda::ndegTwistedMassGPU";
24 template <
typename Dslash>
27 static_assert(
xpay ==
true,
"Non-generate twisted-mass operator only defined for xpay");
28 dslash.
launch(ndegTwistedMassGPU<Float, nDim, nColor, nParity, dagger, xpay, kernel_type, Arg>, tp, arg, stream);
41 Dslash<Float>(arg, out, in,
"kernels/dslash_ndeg_twisted_mass.cuh"),
57 errorQuda(
"Non-degenerate twisted-mass operator only defined for xpay=true");
63 switch (arg.kernel_type) {
89 constexpr
int nDim = 4;
90 NdegTwistedMassArg<Float, nColor, recon> arg(out, in, U, a, b, c, x, parity, dagger, comm_override);
94 twisted, const_cast<cudaColorSpinorField *>(static_cast<const cudaColorSpinorField *>(&in)),
106 #ifdef GPU_NDEG_TWISTED_MASS_DIRAC 107 if (in.
V() == out.
V())
errorQuda(
"Aliasing pointers");
117 instantiate<NdegTwistedMassApply>(
out,
in, U, a, b, c, x,
parity,
dagger, comm_override, profile);
119 errorQuda(
"Non-degenerate twisted-mass dslash has not been built");
120 #endif // GPU_NDEG_TWISTED_MASS_DIRAC void launch(T *f, const TuneParam &tp, Arg &arg, const cudaStream_t &stream)
void apply(const cudaStream_t &stream)
QudaVerbosity getVerbosity()
#define checkPrecision(...)
const char * VolString() const
static void launch(Dslash &dslash, TuneParam &tp, Arg &arg, const cudaStream_t &stream)
static constexpr const char * kernel
void xpay(ColorSpinorField &x, double a, ColorSpinorField &y)
virtual ~NdegTwistedMass()
void ApplyNdegTwistedMass(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b, double c, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Driver for applying the non-degenerate twisted-mass stencil.
TuneParam & tuneLaunch(Tunable &tunable, QudaTune enabled, QudaVerbosity verbosity)
#define checkLocation(...)
int ghostFaceCB[QUDA_MAX_DIM+1]
const DslashConstant & getDslashConstant() const
Get the dslash_constant structure from this field.
cpuColorSpinorField * out
void apply(const cudaStream_t &stream)
void resizeVector(int y, int z) const
This is a helper class that is used to instantiate the correct templated kernel for the dslash...
const ColorSpinorField & in
__host__ __device__ ValueType arg(const complex< ValueType > &z)
Returns the phase angle of z.
virtual long long flops() const
NdegTwistedMassApply(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b, double c, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
QudaTune getTuning()
Query whether autotuning is enabled or not. Default is enabled but can be overridden by setting QUDA_...
NdegTwistedMass(Arg &arg, const ColorSpinorField &out, const ColorSpinorField &in)
QudaFieldOrder FieldOrder() const