Kernels Namespace Reference

CUDA Kernels. More...

Functions
__global__ void	Average_Plaquette (float *hgs_d, float *hgt_d, Complex_f *u11t, Complex_f *u12t, unsigned int *iu)
	Calculates the gauge action using new (how new?) lookup table Follows a routine called qedplaq in some QED3 code.
__global__ void	Polyakov (Complex_f *Sigma11, Complex_f *Sigma12, Complex_f *u11t, Complex_f *u12t)
	Calculate the Polyakov loop (no prizes for guessing that one...).
template<typename T>
__global__ void	Half_Leaves (complex< T > *hLeaves0, complex< T > *hLeaves1, complex< T > *u11t, complex< T > *u12t, unsigned int *iu, unsigned int *id, const unsigned short mu, const unsigned short nu)
	Calculates the products of the first two links in a plaquette.
template<typename T>
__global__ void	Full_Clover (complex< T > *clover1, complex< T > *clover2, complex< T > *u11t, complex< T > *u12t, unsigned int *iu, unsigned int *id, int mu, int nu)
	Calculates the clovers in all directions at all sites \( F_{\mu\nu}(n)=\frac{-i}{8a^2}\left(Q_{\mu\nu}(n)-Q_{\nu\mu}(n)\right)\).
template<typename T>
__global__ void	cuCalcXmunu (Bilinear_a Xmunu, const complex< T > *X1, const complex< T > *X2, const complex< T > *sigval, const unsigned short *sigin, const unsigned short clov)
	Gets \(X_{\mu\nu}\) for the clover force.
template<typename T>
__global__ void	Clov_Force (double *dSdpi, const complex< T > *u11t, const complex< T > *u12t, Bilinear_a Xmn, const complex< T > *sigval, const unsigned short *sigin, const unsigned int *iu, const unsigned int *id, const float akappa, const unsigned short mu, const unsigned short nu)
	cuGets the clover contribution to the force
template<typename T>
__global__ void	ByClover (complex< T > *phi, complex< T > *r, complex< T > *clover1, complex< T > *clover2, complex< T > *sigval, const float akappa, unsigned short *sigin, bool dag)
	Clover analogue of the Dslash operation. This version acts on all flavours similar to Dslash and Dslash_d.
template<typename T>
__global__ void	HbyClover (complex< T > *phi, complex< T > *r, complex< T > *clover1, complex< T > *clover2, complex< T > *sigval, const float akappa, unsigned short *sigin, bool dag)
	Clover analogue of the Hdslash operation. The H in front is for half, as we only act on the fermions of flavour 1.
__global__ void	Plus_staple (const int mu, const int nu, unsigned int *iu, Complex_f *Sigma11, Complex_f *Sigma12, Complex_f *u11t, Complex_f *u12t)
	Calculates the staple in the positive \(\mu\) direction.
__global__ void	Minus_staple (const int mu, const int nu, unsigned int *iu, unsigned int *id, Complex_f *Sigma11, Complex_f *Sigma12, Complex_f *u11sh, Complex_f *u12sh, Complex_f *u11t, Complex_f *u12t)
	Calculates the staple in the negative \(\mu\) direction.
__global__ void	cuGaugeForce (int mu, Complex_f *Sigma11, Complex_f *Sigma12, double *dSdpi, Complex_f *u11t, Complex_f *u12t, float beta)
	Calculates the gauge force due to the Wilson Action at each intermediate time.
template<typename T>
__global__ void	Gather (T *x, T *y, const unsigned int n, unsigned int *table, const unsigned short mu)
	Extracts all the single precision gauge links in the \(\mu\) direction only.
__global__ void	cuForce_s (double *dSdpi, Complex_f *u11t, Complex_f *u12t, Complex_f *X1, Complex_f *X2, Complex_f gamval[20], unsigned int *iu, const unsigned short gamin[16], float akappa, const unsigned short mu)
	Calculates the force \(\frac{dS}{d\pi}\) at each intermediate time.
__global__ void	cuForce_t (double *dSdpi, Complex_f *u11t, Complex_f *u12t, Complex_f *X1, Complex_f *X2, Complex_f gamval[20], float *dk4m, float *dk4p, unsigned int *iu, const unsigned short gamin[16], float akappa)
	Calculates the force \(\frac{dS}{d\pi}\) at each intermediate time.
template<typename T>
__global__ void	cuDslash (complex< T > *phi, complex< T > *r, complex< T > *u11t, complex< T > *u12t, const unsigned int *iu, const unsigned int *id, complex< T > gamval[20], const unsigned short gamin[16], const T *dk4m, const T *dk4p, const Complex_f jqq, const float akappa)
	Evaluates \(\Phi=M r\).
template<typename T>
__global__ void	cuDslashd (complex< T > *phi, const complex< T > *r, const complex< T > *u11t, const complex< T > *u12t, const unsigned int *iu, const unsigned int *id, complex< T > gamval[20], const unsigned short gamin[16], const T *dk4m, const T *dk4p, const Complex_f jqq, const float akappa)
	Evaluates \(\Phi=M^\dagger r\).
template<typename T>
__global__ void	cuHdslash (complex< T > *phi, const complex< T > *r, const complex< T > *u11t, const complex< T > *u12t, unsigned int *iu, unsigned int *id, __constant__ complex< T > gamval[20], const unsigned short gamin[16], const T *dk4m, const T *dk4p, const __grid_constant__ float akappa)
	Evaluates \(\Phi=Mr\) using up/down partitioning.
template<typename T>
__global__ void	cuHdslashd (complex< T > *phi, const complex< T > *r, const complex< T > *u11t, const complex< T > *u12t, unsigned int *iu, unsigned int *id, __constant__ complex< T > gamval[20], const unsigned short gamin[16], const T *dk4m, const T *dk4p, const __grid_constant__ float akappa)
	Evaluates \(\Phi=M^\dagger r\) using up/down partitioning.
template<typename T>
__global__ void	Transpose (T *out, const T *in, const int fast_in, const int fast_out)
	Swaps the order of the gauge field so that it is now SoA instead of AoS and it is nice and coalesced in memory.
__global__ void	Mixed_Sumto (double *d, float *f, const unsigned int n)
	Sums a float array into a double array.
template<typename T, unsigned int bsize>
__global__ void	reduce_sum (T *g_in_data, T *g_out_data, const unsigned int n)
	Performs a block reduction for sum.
__global__ void	Real_convert (float *a, double *b, const unsigned int len, const bool dtof)
	takes an array of real float and double precision numbers and converts the precision
__global__ void	cuFill_Small_Phi (const unsigned int na, Complex *smallPhi, Complex *Phi)
	Copies necessary (2*4*kvol) elements of Phi into a vector variable.
__global__ void	cuUpDownPart (const unsigned int na, Complex *X0, Complex *R1)
	Up/Down partitioning of the pseudofermion field.
template<typename T>
__global__ void	cuReunitarise (complex< T > *u11t, complex< T > *u12t)
	Reunitarises u11t and u12t as in conj(u11t[i])*u11t[i]+conj(u12t[i])*u12t[i]=1.
__global__ void	cuGauge_Update (const double d, double *pp, Complex *u11t, Complex *u12t, int mu)
	Gauge update for the integration step of the HMC.

Detailed Description

CUDA Kernels.