cusu2hmc.cu File Reference

Several random but useful CUDA routines. More...

#include <cuda_runtime.h>
#include <su2hmc.h>

Include dependency graph for cusu2hmc.cu:

Go to the source code of this file.

Namespaces
namespace	Kernels
	CUDA Kernels.

Macros
#define	MIN(x, y)
	Return the max of two values.
#define	MAX(x, y)
	Return the min of two values.

Functions
template<typename T>
__device__ __forceinline__ T	conj (const T &z)
	Complex Conjugation.
__global__ void	Kernels::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	Kernels::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	Kernels::cuUpDownPart (const unsigned int na, Complex *X0, Complex *R1)
	Up/Down partitioning of the pseudofermion field.
template<typename T>
__global__ void	Kernels::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	Kernels::cuGauge_Update (const double d, double *pp, Complex *u11t, Complex *u12t, int mu)
	Gauge update for the integration step of the HMC.
void	blockInit (int x, int y, int z, int t, dim3 *dimBlock, dim3 *dimGrid)
	Initialises the CUDA grid and block size for a given lattice.
void	Init_CUDA (Complex *u11t, Complex *u12t, Complex gamval[20], Complex_f gamval_f[20], unsigned short gamin[16], double *dk4m, double *dk4p, unsigned int *iu, unsigned int *id)
	Initialise CUDA cuInit was taken already by CUDA (unsurprisingly).
void	cuReal_convert (float *a, double *b, const unsigned int len, const bool dtof, dim3 dimBlock, dim3 dimGrid)
	takes an array of real-valued float and double precision numbers and converts the precision
void	cuComplex_convert (Complex_f *a, Complex *b, const unsigned int len, const bool dtof, dim3 dimBlock, dim3 dimGrid)
	takes an array of complex float and double precision numbers and converts the precision
void	cuFill_Small_Phi (const unsigned int na, Complex *smallPhi, Complex *Phi, dim3 dimBlock, dim3 dimGrid)
	Copies necessary (2*4*kvol) elements of Phi into a vector variable.
void	cuUpDownPart (const unsigned int na, Complex *X0, Complex *R1, dim3 dimBlock, dim3 dimGrid)
	Up/Down partitioning of the pseudofermion field.
void	cuReunitarise (Complex *ut[2], dim3 dimGrid, dim3 dimBlock)
	Reunitarises u11t and u12t as in conj(u11t[i])*u11t[i]+conj(u12t[i])*u12t[i]=1.
void	cuGauge_Update (const double d, double *pp, Complex *ut[2], dim3 dimGrid, dim3 dimBlock)
	CUDA wrapper for the gauge update during the integration step of the HMC.
template __device__ __forceinline__ complex< float >	conj (const complex< float > &z)
template __device__ __forceinline__ complex< double >	conj (const complex< double > &z)

Variables
dim3	dimBlockOne = dim3(1,1,1)
	block size of one
dim3	dimGridOne = dim3(1,1,1)
	Grid size of one.
dim3	dimBlock = dim3(1,1,1)
	Default block size. Usually 128.
dim3	dimGrid = dim3(1,1,1)
	Default grid size. First component is normally nt. Second and third depend whatever is needed to get the block size to 128.
cudaStream_t	streams [ndirac *ndim *nadj]
	An array of concurrent GPU streams to keep it busy.

Detailed Description

Several random but useful CUDA routines.

Author

D. Lawlor

Definition in file cusu2hmc.cu.

Variable Documentation

dimBlock

dim3 dimBlock = dim3(1,1,1)

Default block size. Usually 128.

Grids and Blocks

Threads are grouped together to form warps of 32 threads best to keep the block dimension (ksizex*ksizey) multiples of 32, usually between 128 and 256 Note that from Volta/Turing each SM (group of processors) is smaller than on previous and subsequent generations of GPUs. This means less threads accessing the same number of registers which can increase occupancy for register heavy kernels.

Definition at line 25 of file cusu2hmc.cu.

dimBlockOne

dim3 dimBlockOne = dim3(1,1,1)

block size of one

Definition at line 22 of file cusu2hmc.cu.

dimGrid

dim3 dimGrid = dim3(1,1,1)

Default grid size. First component is normally nt. Second and third depend whatever is needed to get the block size to 128.

Definition at line 27 of file cusu2hmc.cu.

dimGridOne

dim3 dimGridOne = dim3(1,1,1)

Grid size of one.

Definition at line 23 of file cusu2hmc.cu.

streams

cudaStream_t streams[ndirac *ndim *nadj]

An array of concurrent GPU streams to keep it busy.

Definition at line 29 of file cusu2hmc.cu.