Inheritance diagram for quda::Gamma< ValueType, basis, dir >:

Collaboration diagram for quda::Gamma< ValueType, basis, dir >:

[legend]

Public Member Functions
__device__ __host__	Gamma ()

__device__ __host__	Gamma (const Gamma &g)

__device__ __host__ int	getcol (int row) const

__device__ __host__ complex< ValueType >	getelem (int row) const

__device__ __host__ complex< ValueType >	getelem (int row, int col) const

__device__ __host__ complex< ValueType >	getrowelem (int row, int &col) const

__device__ __host__ complex< ValueType >	apply (int row, const complex< ValueType > &a) const

constexpr int	Dir () const

	Gamma (Arg &arg, const ColorSpinorField &meta)

virtual	~Gamma ()

void	apply (const cudaStream_t &stream)

TuneKey	tuneKey () const

void	preTune ()

void	postTune ()

Public Member Functions inherited from quda::TunableVectorY
	TunableVectorY (unsigned int vector_length_y)

bool	advanceBlockDim (TuneParam &param) const

void	initTuneParam (TuneParam &param) const

void	defaultTuneParam (TuneParam &param) const

void	resizeVector (int y) const

void	resizeStep (int y) const

Public Member Functions inherited from quda::Tunable
	Tunable ()

virtual	~Tunable ()

virtual int	tuningIter () const

virtual std::string	paramString (const TuneParam &param) const

virtual std::string	perfString (float time) const

virtual bool	advanceTuneParam (TuneParam &param) const

void	checkLaunchParam (TuneParam &param)

CUresult	jitifyError () const

CUresult &	jitifyError ()

Protected Member Functions
long long	flops () const

long long	bytes () const

bool	tuneGridDim () const

unsigned int	minThreads () const

Protected Member Functions inherited from quda::TunableVectorY
virtual unsigned int	sharedBytesPerThread () const

virtual unsigned int	sharedBytesPerBlock (const TuneParam &param) const

Protected Member Functions inherited from quda::Tunable
virtual bool	tuneAuxDim () const

virtual bool	tuneSharedBytes () const

virtual bool	advanceGridDim (TuneParam &param) const

virtual unsigned int	maxBlockSize (const TuneParam &param) const

virtual unsigned int	maxGridSize () const

virtual unsigned int	minGridSize () const

virtual int	gridStep () const
	gridStep sets the step size when iterating the grid size in advanceGridDim. More...

virtual int	blockStep () const

virtual int	blockMin () const

virtual void	resetBlockDim (TuneParam &param) const

unsigned int	maxBlocksPerSM () const
	For some reason this can't be queried from the device properties, so here we set set this. Based on Table 14 of the CUDA Programming Guide 10.0 (Technical Specifications per Compute Capability) More...

template<typename F >
void	setMaxDynamicSharedBytesPerBlock (F *func) const
	Enable the maximum dynamic shared bytes for the kernel "func" (values given by maxDynamicSharedBytesPerBlock()). More...

unsigned int	maxDynamicSharedBytesPerBlock () const
	This can't be correctly queried in CUDA for all architectures so here we set set this. Based on Table 14 of the CUDA Programming Guide 10.0 (Technical Specifications per Compute Capability). More...

virtual unsigned int	maxSharedBytesPerBlock () const
	The maximum shared memory that a CUDA thread block can use in the autotuner. This isn't necessarily the same as maxDynamicSharedMemoryPerBlock since that may need explicit opt in to enable (by calling setMaxDynamicSharedBytes for the kernel in question). If the CUDA kernel in question does this opt in then this function can be overloaded to return maxDynamicSharedBytesPerBlock. More...

virtual bool	advanceSharedBytes (TuneParam &param) const

virtual bool	advanceAux (TuneParam &param) const

int	writeAuxString (const char *format,...)

Protected Attributes
Arg &	arg

const ColorSpinorField &	meta

Protected Attributes inherited from quda::TunableVectorY
unsigned int	vector_length_y

unsigned int	step_y

bool	tune_block_x

Protected Attributes inherited from quda::Tunable
char	aux [TuneKey::aux_n]

CUresult	jitify_error

Static Private Attributes
static constexpr int	ndim = 4

Detailed Description

template<typename ValueType, QudaGammaBasis basis, int dir>
class quda::Gamma< ValueType, basis, dir >

Definition at line 9 of file gamma.cuh.

Constructor & Destructor Documentation

◆ Gamma() [1/3]

template<typename ValueType, QudaGammaBasis basis, int dir>

__device__ __host__ quda::Gamma< ValueType, basis, dir >::Gamma ( )

inline

Definition at line 25 of file gamma.cuh.

◆ Gamma() [2/3]

template<typename ValueType, QudaGammaBasis basis, int dir>

__device__ __host__ quda::Gamma< ValueType, basis, dir >::Gamma ( const Gamma< ValueType, basis, dir > & g )

inline

Definition at line 26 of file gamma.cuh.

◆ Gamma() [3/3]

template<typename ValueType, QudaGammaBasis basis, int dir>

quda::Gamma< ValueType, basis, dir >::Gamma	(	Arg &	arg,
		const ColorSpinorField &	meta
	)

inline

Definition at line 266 of file dslash_quda.cu.

References quda::LatticeField::AuxString().

Here is the call graph for this function:

◆ ~Gamma()

template<typename ValueType, QudaGammaBasis basis, int dir>

virtual quda::Gamma< ValueType, basis, dir >::~Gamma ( )

inlinevirtual

Definition at line 270 of file dslash_quda.cu.

Member Function Documentation

◆ apply() [1/2]

template<typename ValueType, QudaGammaBasis basis, int dir>

__device__ __host__ complex<ValueType> quda::Gamma< ValueType, basis, dir >::apply	(	int	row,
		const complex< ValueType > &	a
	)		const

inline

Definition at line 221 of file gamma.cuh.

References QUDA_DEGRAND_ROSSI_GAMMA_BASIS, and QUDA_UKQCD_GAMMA_BASIS.

Referenced by quda::ApplyGamma(), and quda::multiplyVUV().

Here is the caller graph for this function:

◆ apply() [2/2]

template<typename ValueType, QudaGammaBasis basis, int dir>

void quda::Gamma< ValueType, basis, dir >::apply ( const cudaStream_t & stream )

inlinevirtual

Implements quda::Tunable.

Definition at line 272 of file dslash_quda.cu.

References quda::arg(), quda::TuneParam::block, errorQuda, getTuning(), getVerbosity(), quda::TuneParam::grid, quda::LatticeField::Location(), QUDA_CPU_FIELD_LOCATION, quda::TuneParam::shared_bytes, and quda::tuneLaunch().

Here is the call graph for this function:

◆ bytes()

template<typename ValueType, QudaGammaBasis basis, int dir>

long long quda::Gamma< ValueType, basis, dir >::bytes ( ) const

inlineprotectedvirtual

Reimplemented from quda::Tunable.

Definition at line 261 of file dslash_quda.cu.

◆ Dir()

template<typename ValueType, QudaGammaBasis basis, int dir>

constexpr int quda::Gamma< ValueType, basis, dir >::Dir ( ) const

inline

Definition at line 292 of file gamma.cuh.

◆ flops()

template<typename ValueType, QudaGammaBasis basis, int dir>

long long quda::Gamma< ValueType, basis, dir >::flops ( ) const

inlineprotectedvirtual

Implements quda::Tunable.

Definition at line 260 of file dslash_quda.cu.

◆ getcol()

template<typename ValueType, QudaGammaBasis basis, int dir>

__device__ __host__ int quda::Gamma< ValueType, basis, dir >::getcol ( int row ) const

inline

Definition at line 28 of file gamma.cuh.

References QUDA_DEGRAND_ROSSI_GAMMA_BASIS.

Referenced by quda::Gamma< ValueType, basis, dir >::getelem(), quda::Gamma< ValueType, basis, dir >::getrowelem(), and quda::multiplyVUV().

Here is the caller graph for this function:

◆ getelem() [1/2]

template<typename ValueType, QudaGammaBasis basis, int dir>

__device__ __host__ complex<ValueType> quda::Gamma< ValueType, basis, dir >::getelem ( int row ) const

inline

Definition at line 98 of file gamma.cuh.

References QUDA_DEGRAND_ROSSI_GAMMA_BASIS, and QUDA_UKQCD_GAMMA_BASIS.

Referenced by quda::Gamma< ValueType, basis, dir >::getelem(), and quda::Gamma< ValueType, basis, dir >::getrowelem().

Here is the caller graph for this function:

◆ getelem() [2/2]

template<typename ValueType, QudaGammaBasis basis, int dir>

__device__ __host__ complex<ValueType> quda::Gamma< ValueType, basis, dir >::getelem	(	int	row,
		int	col
	)		const

inline

Definition at line 209 of file gamma.cuh.

References quda::Gamma< ValueType, basis, dir >::getcol(), and quda::Gamma< ValueType, basis, dir >::getelem().

Here is the call graph for this function:

◆ getrowelem()

template<typename ValueType, QudaGammaBasis basis, int dir>

__device__ __host__ complex<ValueType> quda::Gamma< ValueType, basis, dir >::getrowelem	(	int	row,
		int &	col
	)		const

inline

Definition at line 215 of file gamma.cuh.

References quda::Gamma< ValueType, basis, dir >::getcol(), and quda::Gamma< ValueType, basis, dir >::getelem().

Here is the call graph for this function:

◆ minThreads()

template<typename ValueType, QudaGammaBasis basis, int dir>

unsigned int quda::Gamma< ValueType, basis, dir >::minThreads ( ) const

inlineprotectedvirtual

Reimplemented from quda::Tunable.

Definition at line 263 of file dslash_quda.cu.

References quda::Arg< real, Ns, Nc, order >::volumeCB.

◆ postTune()

template<typename ValueType, QudaGammaBasis basis, int dir>

void quda::Gamma< ValueType, basis, dir >::postTune ( )

inlinevirtual

Reimplemented from quda::Tunable.

Definition at line 287 of file dslash_quda.cu.

◆ preTune()

template<typename ValueType, QudaGammaBasis basis, int dir>

void quda::Gamma< ValueType, basis, dir >::preTune ( )

inlinevirtual

Reimplemented from quda::Tunable.

Definition at line 286 of file dslash_quda.cu.

◆ tuneGridDim()

template<typename ValueType, QudaGammaBasis basis, int dir>

bool quda::Gamma< ValueType, basis, dir >::tuneGridDim ( ) const

inlineprotectedvirtual

Reimplemented from quda::Tunable.

Definition at line 262 of file dslash_quda.cu.

◆ tuneKey()

template<typename ValueType, QudaGammaBasis basis, int dir>

TuneKey quda::Gamma< ValueType, basis, dir >::tuneKey ( ) const

inlinevirtual

Implements quda::Tunable.

Definition at line 284 of file dslash_quda.cu.

References quda::LatticeField::VolString().

Here is the call graph for this function:

Member Data Documentation

◆ arg

template<typename ValueType, QudaGammaBasis basis, int dir>

Arg& quda::Gamma< ValueType, basis, dir >::arg

protected

Definition at line 257 of file dslash_quda.cu.

◆ meta

template<typename ValueType, QudaGammaBasis basis, int dir>

const ColorSpinorField& quda::Gamma< ValueType, basis, dir >::meta

protected

Definition at line 258 of file dslash_quda.cu.

◆ ndim

template<typename ValueType, QudaGammaBasis basis, int dir>

constexpr int quda::Gamma< ValueType, basis, dir >::ndim = 4

staticprivate

Definition at line 11 of file gamma.cuh.

The documentation for this class was generated from the following files:

include/gamma.cuh
lib/dslash_quda.cu

Public Member Functions

Protected Member Functions

Protected Attributes

Static Private Attributes

Detailed Description

template<typename ValueType, QudaGammaBasis basis, int dir> class quda::Gamma< ValueType, basis, dir >

Constructor & Destructor Documentation

◆ Gamma() [1/3]

◆ Gamma() [2/3]

◆ Gamma() [3/3]

◆ ~Gamma()

Member Function Documentation

◆ apply() [1/2]

◆ apply() [2/2]

◆ bytes()

◆ Dir()

◆ flops()

◆ getcol()

◆ getelem() [1/2]

◆ getelem() [2/2]

◆ getrowelem()

◆ minThreads()

◆ postTune()

◆ preTune()

◆ tuneGridDim()

◆ tuneKey()

Member Data Documentation

◆ arg

◆ meta

◆ ndim

template<typename ValueType, QudaGammaBasis basis, int dir>
class quda::Gamma< ValueType, basis, dir >