20 template <
typename Float,
int nDim,
int nColor,
int nParity,
bool dagger,
bool xpay, KernelType kernel_type,
typename Arg>
22 static constexpr
const char *
kernel =
"quda::twistedMassPreconditionedGPU";
23 template <
typename Dslash>
26 static_assert(nParity == 1,
"preconditioned twisted-mass operator only defined for nParity=1");
28 twistedMassPreconditionedGPU<Float, nDim, nColor, nParity, dagger, xpay, kernel_type, Arg>, tp, arg, stream);
41 Dslash<Float>(arg, out, in,
"kernels/dslash_twisted_mass_preconditioned.cuh"),
46 for (
int i = 0; i < 8; i++)
56 if (arg.asymmetric && !arg.dagger)
errorQuda(
"asymmetric operator only defined for dagger");
57 if (arg.asymmetric && arg.xpay)
errorQuda(
"asymmetric operator not defined for xpay");
65 errorQuda(
"Preconditioned twisted-mass operator not defined nParity=%d", arg.
nParity);
72 switch (arg.kernel_type) {
98 constexpr
int nDim = 4;
99 TwistedMassArg<Float, nColor, recon> arg(out, in, U, a, b, xpay, x, parity, dagger, asymmetric, comm_override);
103 const_cast<cudaColorSpinorField *>(static_cast<const cudaColorSpinorField *>(&in)), in.
VolumeCB(),
120 #ifdef GPU_TWISTED_MASS_DIRAC 121 if (in.
V() == out.
V())
errorQuda(
"Aliasing pointers");
134 instantiate<TwistedMassPreconditionedApply>(
135 out,
in, U, a, b,
xpay, x,
parity,
dagger, asymmetric, comm_override, profile);
139 errorQuda(
"Twisted-mass dslash has not been built");
140 #endif // GPU_TWISTED_MASS_DIRAC TwistedMassPreconditionedApply(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b, bool xpay, const ColorSpinorField &x, int parity, bool dagger, bool asymmetric, const int *comm_override, TimeProfile &profile)
void launch(T *f, const TuneParam &tp, Arg &arg, const cudaStream_t &stream)
void apply(const cudaStream_t &stream)
QudaVerbosity getVerbosity()
#define checkPrecision(...)
virtual ~TwistedMassPreconditioned()
const char * VolString() const
TwistedMassPreconditioned(Arg &arg, const ColorSpinorField &out, const ColorSpinorField &in)
void xpay(ColorSpinorField &x, double a, ColorSpinorField &y)
static void launch(Dslash &dslash, TuneParam &tp, Arg &arg, const cudaStream_t &stream)
const int * GhostFaceCB() const
TuneParam & tuneLaunch(Tunable &tunable, QudaTune enabled, QudaVerbosity verbosity)
#define checkLocation(...)
const ColorSpinorField & in
static constexpr const char * kernel
cpuColorSpinorField * out
__host__ __device__ ValueType arg(const complex< ValueType > &z)
Returns the phase angle of z.
void ApplyTwistedMassPreconditioned(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, double b, bool xpay, const ColorSpinorField &x, int parity, bool dagger, bool asymmetric, const int *comm_override, TimeProfile &profile)
Driver for applying the preconditioned twisted-mass stencil.
void pushKernelPackT(bool pack)
This is a helper class that is used to instantiate the correct templated kernel for the dslash...
virtual long long flops() const
QudaTune getTuning()
Query whether autotuning is enabled or not. Default is enabled but can be overridden by setting QUDA_...
void apply(const cudaStream_t &stream)
QudaFieldOrder FieldOrder() const