21 template <
typename Float,
int nDim,
int nColor,
int nParity,
bool dagger,
bool xpay, KernelType kernel_type,
typename Arg>
22 struct StaggeredLaunch {
23 static constexpr
const char *
kernel =
"quda::staggeredGPU";
24 template <
typename Dslash>
27 dslash.
launch(staggeredGPU<Float, nDim, nColor, nParity, dagger, xpay, kernel_type, Arg>, tp, arg, stream);
31 template <
typename Float,
int nDim,
int nColor,
typename Arg>
class Staggered :
public Dslash<Float>
40 Dslash<Float>(arg, out, in,
"kernels/dslash_staggered.cuh"),
51 errorQuda(
"Staggered Dslash not implemented on CPU");
65 template <
typename Float,
int nColor, QudaReconstructType recon_u>
struct StaggeredApply {
73 #ifdef BUILD_MILC_INTERFACE 74 constexpr
int nDim = 4;
75 constexpr
bool improved =
false;
78 out, in, U, U, a, x, parity, dagger, comm_override);
82 staggered, const_cast<cudaColorSpinorField *>(static_cast<const cudaColorSpinorField *>(&in)), in.
VolumeCB(),
86 errorQuda(
"MILC interface has not been built so MILC phase staggered fermions not enabled");
89 #ifdef BUILD_TIFR_INTERFACE 90 constexpr
int nDim = 4;
91 constexpr
bool improved =
false;
94 out, in, U, U, a, x, parity, dagger, comm_override);
98 staggered, const_cast<cudaColorSpinorField *>(static_cast<const cudaColorSpinorField *>(&in)), in.
VolumeCB(),
102 errorQuda(
"TIFR interface has not been built so TIFR phase taggered fermions not enabled");
116 #ifdef GPU_STAGGERED_DIRAC 117 if (in.
V() == out.
V())
errorQuda(
"Aliasing pointers");
127 instantiate<StaggeredApply, StaggeredReconstruct>(
out,
in, U, a, x,
parity,
dagger, comm_override, profile);
129 errorQuda(
"Staggered dslash has not been built");
void launch(T *f, const TuneParam &tp, Arg &arg, const cudaStream_t &stream)
void apply(const cudaStream_t &stream)
QudaVerbosity getVerbosity()
#define checkPrecision(...)
static void launch(Dslash &dslash, TuneParam &tp, Arg &arg, const cudaStream_t &stream)
const char * VolString() const
StaggeredApply(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
__device__ __host__ void staggered(Arg &arg, int idx, int parity)
const int * GhostFaceCB() const
TuneParam & tuneLaunch(Tunable &tunable, QudaTune enabled, QudaVerbosity verbosity)
#define checkLocation(...)
Main header file for host and device accessors to GaugeFields.
QudaFieldLocation Location() const
Parameter structure for driving the Staggered Dslash operator.
cpuColorSpinorField * out
static constexpr const char * kernel
void ApplyStaggered(ColorSpinorField &out, const ColorSpinorField &in, const GaugeField &U, double a, const ColorSpinorField &x, int parity, bool dagger, const int *comm_override, TimeProfile &profile)
Apply the staggered dslash operator to a color-spinor field.
void apply(const cudaStream_t &stream)
__host__ __device__ ValueType arg(const complex< ValueType > &z)
Returns the phase angle of z.
Staggered(Arg &arg, const ColorSpinorField &out, const ColorSpinorField &in)
QudaStaggeredPhase StaggeredPhase() const
QudaTune getTuning()
Query whether autotuning is enabled or not. Default is enabled but can be overridden by setting QUDA_...
QudaFieldOrder FieldOrder() const