cumatrices.cu File Reference

CUDA version of matrix manipulation routines. More...

#include <assert.h>
#include <su2hmc.h>
#include <matrices.h>
#include <thrust_complex.h>

Include dependency graph for cumatrices.cu:

Go to the source code of this file.

Namespaces
namespace	Device
	CUDA Device code.
namespace	Kernels
	CUDA Kernels.

Macros
#define	SU2_WARP_MASK   0xffffffffU

Functions
template<typename T, unsigned int bsize>
__device__ void	Device::warpReduce_sum (volatile T *sdata, const unsigned int tid)
	Performs a warp reduction for sum.
template<typename T>
__global__ void	Kernels::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	Kernels::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	Kernels::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	Kernels::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	Kernels::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	Kernels::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	Kernels::reduce_sum (T *g_in_data, T *g_out_data, const unsigned int n)
	Performs a block reduction for sum.
double	cureduce_sum_d (double *input, const unsigned int n, const unsigned short stream)
	Sum all terms in an array of doubles.
void	cuDslash (Complex *phi, Complex *r, Complex *ut[nc], unsigned int *iu, unsigned int *id, Complex gamval[20], const unsigned short gamin[16], double *dk[nc], Complex_f jqq, float akappa, dim3 dimGrid, dim3 dimBlock)
	GPU calling wrapper for \(\Phi=M r\) in double precision.
void	cuDslashd (Complex *phi, Complex *r, Complex *ut[nc], unsigned int *iu, unsigned int *id, Complex gamval[20], const unsigned short gamin[16], double *dk[nc], Complex_f jqq, float akappa, dim3 dimGrid, dim3 dimBlock)
	GPU calling wrapper for \(\Phi=M^\dagger r\) in double precision.
void	cuHdslash (Complex *phi, Complex *r, Complex *ut[nc], unsigned int *iu, unsigned int *id, Complex gamval[20], const unsigned short gamin[16], double *dk[nc], float akappa, dim3 dimGrid, dim3 dimBlock)
	GPU calling wrapper for \(\Phi=M^\dagger r\) in double precision.
void	cuHdslashd (Complex *phi, Complex *r, Complex *ut[nc], unsigned int *iu, unsigned int *id, Complex gamval[20], const unsigned short gamin[16], double *dk[nc], float akappa, dim3 dimGrid, dim3 dimBlock)
	GPU calling wrapper for \(\Phi=M^\dagger r\) in double precision.
void	cuDslash_f (Complex_f *phi, Complex_f *r, Complex_f *ut[nc], unsigned int *iu, unsigned int *id, Complex_f gamval[20], const unsigned short gamin[16], float *dk[nc], Complex_f jqq, float akappa, dim3 dimGrid, dim3 dimBlock)
	GPU calling wrapper for \(\Phi=M^\dagger r\) in double precision.
void	cuDslashd_f (Complex_f *phi, Complex_f *r, Complex_f *ut[nc], unsigned int *iu, unsigned int *id, Complex_f gamval[20], const unsigned short gamin[16], float *dk[nc], Complex_f jqq, float akappa, dim3 dimGrid, dim3 dimBlock)
	GPU calling wrapper for \(\Phi=M^\dagger r\) in double precision.
void	cuHdslash_f (Complex_f *phi, Complex_f *r, Complex_f *ut[nc], unsigned int *iu, unsigned int *id, Complex_f gamval[20], const unsigned short gamin[16], float *dk[nc], float akappa, dim3 dimGrid, dim3 dimBlock)
	GPU calling wrapper for \(\Phi=M r\) in single precision.
void	cuHdslashd_f (Complex_f *phi, Complex_f *r, Complex_f *ut[nc], unsigned int *iu, unsigned int *id, Complex_f gamval[20], const unsigned short gamin[16], float *dk[nc], float akappa, dim3 dimGrid, dim3 dimBlock)
	GPU calling wrapper for \(\Phi=M^\dagger r\) in single precision.
void	cuTranspose_z (Complex *out, const int fast_in, const int fast_out, const dim3 dimGrid, const dim3 dimBlock)
	In place transpose used to convert from AoS to SoA memory layout.
void	cuTranspose_c (Complex_f *out, const int fast_in, const int fast_out, const dim3 dimGrid, const dim3 dimBlock)
	In place transpose used to convert from AoS to SoA memory layout.
void	cuTranspose_d (double *out, const int fast_in, const int fast_out, const dim3 dimGrid, const dim3 dimBlock)
	In place transpose used to convert from AoS to SoA memory layout.
void	cuTranspose_f (float *out, const int fast_in, const int fast_out, const dim3 dimGrid, const dim3 dimBlock)
	In place transpose used to convert from AoS to SoA memory layout.
void	cuTranspose_I (int *out, const int fast_in, const int fast_out, const dim3 dimGrid, const dim3 dimBlock)
	In place transpose used to convert from AoS to SoA memory layout.
void	cuTranspose_U (unsigned int *out, const int fast_in, const int fast_out, const dim3 dimGrid, const dim3 dimBlock)
	In place transpose used to convert from AoS to SoA memory layout.
void	cuMixed_Sumto (double *d, float *f, const unsigned int n, const dim3 dimGrid, const dim3 dimBlock)
	Add a single to a double value, and save the output in the double array For complex valued arrays, one may cast the complex<double> and complex<float> arrays to double and float arrays, and use 2N for the array length instead.

Detailed Description

CUDA version of matrix manipulation routines.

Author

D. Lawlor

Definition in file cumatrices.cu.