13 template <
typename real,
typename Arg>
class Contraction : TunableVectorY
17 const ColorSpinorField &x;
18 const ColorSpinorField &y;
22 bool tuneGridDim()
const {
return false; }
23 unsigned int minThreads()
const {
return arg.threads; }
26 Contraction(Arg &arg,
const ColorSpinorField &x,
const ColorSpinorField &y,
const QudaContractType cType) :
36 default:
errorQuda(
"Unexpected contraction type %d", cType);
38 strcat(aux, x.AuxString());
40 create_jitify_program(
"kernels/contraction.cuh");
43 virtual ~Contraction() {}
45 void apply(
const cudaStream_t &
stream)
50 std::string function_name;
54 default:
errorQuda(
"Unexpected contraction type %d", cType);
57 using namespace jitify::reflection;
58 jitify_error = program->kernel(function_name)
59 .instantiate(Type<real>(), Type<Arg>())
60 .configure(tp.grid, tp.block, tp.shared_bytes, stream)
66 computeDegrandRossiContraction<real><<<tp.grid, tp.block, tp.shared_bytes>>>(
arg);
68 default:
errorQuda(
"Unexpected contraction type %d", cType);
76 TuneKey tuneKey()
const {
return TuneKey(x.VolString(),
typeid(*this).name(), aux); }
81 long long flops()
const 84 return 16 * 3 * 6ll * x.Volume();
86 return ((16 * 3 * 6ll) + (16 * (4 + 12))) * x.Volume();
89 long long bytes()
const 91 return x.Bytes() + y.Bytes() + x.Nspin() * x.Nspin() * x.Volume() *
sizeof(complex<real>);
95 template <
typename real>
96 void contract_quda(
const ColorSpinorField &x,
const ColorSpinorField &y, complex<real> *result,
99 ContractionArg<real>
arg(x, y, result);
100 Contraction<real, ContractionArg<real>> contraction(arg, x, y, cType);
101 contraction.apply(0);
126 errorQuda(
"Contraction code has not been built");
QudaVerbosity getVerbosity()
#define checkPrecision(...)
Helper file when using jitify run-time compilation. This file should be included in source code...
QudaGammaBasis GammaBasis() const
#define qudaDeviceSynchronize()
TuneParam & tuneLaunch(Tunable &tunable, QudaTune enabled, QudaVerbosity verbosity)
void contractQuda(const ColorSpinorField &x, const ColorSpinorField &y, void *result, QudaContractType cType)
__host__ __device__ ValueType arg(const complex< ValueType > &z)
Returns the phase angle of z.
enum QudaContractType_s QudaContractType
QudaTune getTuning()
Query whether autotuning is enabled or not. Default is enabled but can be overridden by setting QUDA_...
QudaPrecision Precision() const