Classes
struct	hisq_kernel_param_s

struct	PathCoefficients

struct	RealTypeId< float2 >

struct	RealTypeId< double2 >

struct	CoeffSign

struct	CoeffSign< 0, 1 >

struct	CoeffSign< 0, 0 >

struct	CoeffSign< 1, 1 >

struct	Sign

struct	Sign< 1 >

struct	ArrayLength

struct	ArrayLength< float4 >

class	MiddleLink

class	LepageMiddleLink

class	SideLink

class	SideLinkShort

class	AllLink

class	OneLinkTerm

class	LongLinkTerm

class	CompleteForce

class	DerivativeCoefficients

class	UnitarizeForceCuda

Typedefs
typedef struct quda::fermion_force::hisq_kernel_param_s	hisq_kernel_param_t

Functions
void	hisqForceInitCuda (QudaGaugeParam *param)

void	hisqStaplesForceCuda (const double path_coeff[6], const QudaGaugeParam &param, const cudaGaugeField &oprod, const cudaGaugeField &link, cudaGaugeField *newOprod)

void	hisqLongLinkForceCuda (double coeff, const QudaGaugeParam &param, const cudaGaugeField &oprod, const cudaGaugeField &link, cudaGaugeField *newOprod)

void	hisqCompleteForceCuda (const QudaGaugeParam &param, const cudaGaugeField &oprod, const cudaGaugeField &link, cudaGaugeField *force)

void	setUnitarizeForceConstants (double unitarize_eps, double hisq_force_filter, double max_det_error, bool allow_svd, bool svd_only, double svd_rel_error, double svd_abs_error)

void	unitarizeForceCuda (const QudaGaugeParam &param, cudaGaugeField &cudaOldForce, cudaGaugeField &cudaGauge, cudaGaugeField cudaNewForce, int unitarization_failed)

void	unitarizeForceCPU (const QudaGaugeParam &param, cpuGaugeField &cpuOldForce, cpuGaugeField &cpuGauge, cpuGaugeField *cpuNewForce)

__device__ float2	operator* (float a, const float2 &b)

__device__ double2	operator* (double a, const double2 &b)

__device__ const float2 &	operator+= (float2 &a, const float2 &b)

__device__ const double2 &	operator+= (double2 &a, const double2 &b)

__device__ const float4 &	operator+= (float4 &a, const float4 &b)

template<class T >
__device__ void	adjointMatrix (T *mat)

template<int N, class T >
__device__ void	loadMatrixFromField (const T const field_even, const T const field_odd, int dir, int idx, T *const mat, int oddness, int stride)

template<class T >
__device__ void	loadMatrixFromField (const T const field_even, const T const field_odd, int dir, int idx, T *const mat, int oddness, int stride)

__device__ void	loadMatrixFromField (const float4 const field_even, const float4 const field_odd, int dir, int idx, float2 *const mat, int oddness, int stride)

template<class T >
__device__ void	loadMatrixFromField (const T const field_even, const T const field_odd, int idx, T *const mat, int oddness, int stride)

template<class T , class U >
__device__ void	addMatrixToField (const T const mat, int dir, int idx, U coeff, T const field_even, T *const field_odd, int oddness)

template<class T , class U >
__device__ void	addMatrixToField (const T const mat, int idx, U coeff, T const field_even, T *const field_odd, int oddness)

template<class T , class U >
__device__ void	addMatrixToField_test (const T const mat, int idx, U coeff, T const field_even, T *const field_odd, int oddness)

template<class T >
__device__ void	storeMatrixToField (const T const mat, int dir, int idx, T const field_even, T *const field_odd, int oddness)

template<class T >
__device__ void	storeMatrixToField (const T const mat, int idx, T const field_even, T *const field_odd, int oddness)

template<class T , class U >
__device__ void	storeMatrixToMomentumField (const T const mat, int dir, int idx, U coeff, T const mom_even, T *const mom_odd, int oddness)

template<typename T >
__device__ void	reconstructSign (int *const sign, int dir, const T i[4])

template<class RealA , int oddBit>
__global__ void	do_one_link_term_kernel (const RealA const oprodEven, const RealA const oprodOdd, int sig, typename RealTypeId< RealA >::Type coeff, RealA const outputEven, RealA const outputOdd, const int threads)

template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>
__global__ void	HISQ_KERNEL_NAME (do_middle_link, EXT)(const RealA *const oprodEven

	if (sid >=kparam.threads) return

	if (mu_positive)

	FF_COMPUTE_NEW_FULL_IDX_PLUS_UPDATE (OPP_DIR(mu), new_mem_idx, new_mem_idx)

	if (sig_positive)

	FF_COMPUTE_NEW_FULL_IDX_MINUS_UPDATE (OPP_DIR(sig), new_mem_idx, new_mem_idx)

	if (!mu_positive)

	if (QprevOdd==NULL)

	MATRIX_PRODUCT (bc_link, COLOR_MAT_Y,!mu_positive, COLOR_MAT_W)

	if (PmuOdd)

	MATRIX_PRODUCT (ab_link, COLOR_MAT_W, sig_positive, COLOR_MAT_Y)

	storeMatrixToField (COLOR_MAT_Y, new_sid, P3Even, P3Odd, oddBit)

	adjointMatrix (ad_link)

template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>
__global__ void	HISQ_KERNEL_NAME (do_lepage_middle_link, EXT)(const RealA *const oprodEven

	loadMatrixFromField (oprodEven, oprodOdd, point_c, COLOR_MAT_Y, oddBit, hf.color_matrix_stride)

template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>
__global__ void	HISQ_KERNEL_NAME (do_side_link, EXT)(const RealA *const P3Even

	loadMatrixFromField (P3Even, P3Odd, new_sid, COLOR_MAT_Y, oddBit, hf.color_matrix_stride)

	FF_COMPUTE_NEW_FULL_IDX_PLUS_UPDATE (mymu, new_mem_idx, new_mem_idx)

	RECONSTRUCT_SITE_LINK (ad_link, ad_link_sign)

	MATRIX_PRODUCT (ad_link, COLOR_MAT_Y, mu_positive, COLOR_MAT_W)

	addMatrixToField (COLOR_MAT_W, point_d, accumu_coeff, shortPEven, shortPOdd, 1-oddBit)

	loadMatrixFromField (QprodEven, QprodOdd, point_d, COLOR_MAT_X, 1-oddBit, hf.color_matrix_stride)

	if (oddBit)

	addMatrixToNewOprod (COLOR_MAT_W, OPP_DIR(mu), new_sid, mycoeff, newOprodEven, newOprodOdd, oddBit)

template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>
__global__ void	HISQ_KERNEL_NAME (do_side_link_short, EXT)(const RealA *const P3Even

template<class RealA , class RealB , SHORT sig_positive, SHORT mu_positive, SHORT _oddBit, int oddness_change>
__global__ void	HISQ_KERNEL_NAME (do_all_link, EXT)(const RealA *const oprodEven

	FF_COMPUTE_NEW_FULL_IDX_PLUS_UPDATE (mu, new_mem_idx, new_mem_idx)

	loadMatrixFromField (QprevEven, QprevOdd, point_d, COLOR_MAT_X, 1-oddBit, hf.color_matrix_stride)

	HISQ_LOAD_LINK (linkEven, linkOdd, mu, new_sid, ad_link, oddBit)

	HISQ_LOAD_LINK (linkEven, linkOdd, mu, point_b, bc_link, 1-oddBit)

	RECONSTRUCT_SITE_LINK (bc_link, bc_link_sign)

	MAT_MUL_MAT (bc_link, COLOR_MAT_Y, COLOR_MAT_Z)

	MATRIX_PRODUCT (ab_link, COLOR_MAT_Z, sig_positive, COLOR_MAT_Y)

	ADJ_MAT_MUL_ADJ_MAT (COLOR_MAT_X, COLOR_MAT_Y, COLOR_MAT_W)

	addMatrixToNewOprod (COLOR_MAT_W, mu, new_sid, Sign< _oddBit^oddness_change >::result *mycoeff, newOprodEven, newOprodOdd, oddBit)

	MATRIX_PRODUCT (ad_link, COLOR_MAT_Y, 0, COLOR_MAT_W)

template<class RealA , class RealB , int oddBit>
__global__ void	HISQ_KERNEL_NAME (do_longlink, EXT)(const RealB *const linkEven

	if (GOES_FORWARDS(sig))

template<class RealA , class RealB , int oddBit>
__global__ void	HISQ_KERNEL_NAME (do_complete_force, EXT)(const RealB *const linkEven

	if (sid >=threads) return

	HISQ_LOAD_LINK (linkEven, linkOdd, sig, new_sid, LINK_W, oddBit)

	RECONSTRUCT_SITE_LINK (LINK_W, link_sign)

	loadMatrixFromField (oprodEven, oprodOdd, sig, new_sid, COLOR_MAT_X, oddBit, hf.color_matrix_stride)

	MAT_MUL_MAT (LINK_W, COLOR_MAT_X, COLOR_MAT_W)

	storeMatrixToMomentumField (COLOR_MAT_W, sig, sid, coeff, forceEven, forceOdd, oddBit)

template<class Cmplx >
__device__ __host__ void	accumBothDerivatives (Matrix< Cmplx, 3 > *result, const Matrix< Cmplx, 3 > &left, const Matrix< Cmplx, 3 > &right, const Matrix< Cmplx, 3 > &outer_prod)

template<class Cmplx >
__device__ __host__ void	accumDerivatives (Matrix< Cmplx, 3 > *result, const Matrix< Cmplx, 3 > &left, const Matrix< Cmplx, 3 > &right, const Matrix< Cmplx, 3 > &outer_prod)

template<class T >
__device__ __host__ T	getAbsMin (const T *const array, int size)

template<class Real >
__device__ __host__ bool	checkAbsoluteError (Real a, Real b, Real epsilon)

template<class Real >
__device__ __host__ bool	checkRelativeError (Real a, Real b, Real epsilon)

template<class Cmplx >
__device__ __host__ void	reciprocalRoot (Matrix< Cmplx, 3 > res, DerivativeCoefficients< typename RealTypeId< Cmplx >::Type > deriv_coeffs, typename RealTypeId< Cmplx >::Type f[3], Matrix< Cmplx, 3 > &q, int *unitarization_failed)

template<class Cmplx >
__device__ __host__ void	getUnitarizeForceSite (const Matrix< Cmplx, 3 > &v, const Matrix< Cmplx, 3 > &outer_prod, Matrix< Cmplx, 3 > result, int unitarization_failed)

template<class Cmplx >
__global__ void	getUnitarizeForceField (const int threads, const Cmplx link_even, const Cmplx link_odd, const Cmplx old_force_even, const Cmplx old_force_odd, Cmplx force_even, Cmplx force_odd, int *unitarization_failed)

Variables
texture< int4, 1 >	newOprod0TexDouble

texture< int4, 1 >	newOprod1TexDouble

texture< float2, 1, cudaReadModeElementType >	newOprod0TexSingle

texture< float2, 1, cudaReadModeElementType >	newOprod1TexSingle

__global__ void const RealA *const	oprodOdd

__global__ void const RealA const const RealA const	QprevEven

__global__ void const RealA const const RealA const const RealA *const	QprevOdd

__global__ void const RealA const const RealA const const RealA const const RealB const	linkEven

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB *const	linkOdd

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB *const int	sig

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB *const int int	mu

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB *const int int RealTypeId< RealA >::Type	coeff = (oddBit==1) ? -1 : 1

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const	PmuEven

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const RealA *const	PmuOdd

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const RealA const RealA const	P3Even

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const RealA const RealA const RealA *const	P3Odd

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const RealA const RealA const RealA const RealA const	QmuEven

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const RealA const RealA const RealA const RealA const RealA *const	QmuOdd

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const RealA const RealA const RealA const RealA const RealA const RealA const	newOprodEven

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const RealA const RealA const RealA const RealA const RealA const RealA const RealA *const	newOprodOdd

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealA const RealA const RealA const RealA const RealA const RealA const RealA const RealA *const hisq_kernel_param_t	kparam

int	sid = blockIdx.x * blockDim.x + threadIdx.x

int	x [4] = z1 - z2*D2

int	z1 = sid/D1h

int	x1h = sid - z1*D1h

int	z2 = z1/D2

int	x1odd = (x[1] + x[2] + x[3] + oddBit) & 1

int	new_x [4] = x[0]

int	new_mem_idx = X

RealA	ab_link [ArrayLength< RealA >::result]

RealA	bc_link [ArrayLength< RealA >::result]

RealA	ad_link [ArrayLength< RealA >::result]

RealA	COLOR_MAT_W [ArrayLength< RealA >::result]

RealA	COLOR_MAT_Y [ArrayLength< RealA >::result]

RealA	COLOR_MAT_X [ArrayLength< RealA >::result]

int	point_b = (new_mem_idx >> 1)

int	point_c = (new_mem_idx >> 1)

int	point_d = (new_mem_idx >> 1)

int	ad_link_nbr_idx

int	ab_link_nbr_idx = (sig_positive) ? new_sid : point_b

int	bc_link_nbr_idx

int	mymu

int	X = 2*sid + x1odd

int	new_sid = sid

	else

COMPUTE_LINK_SIGN &	ad_link_sign

int	mysig

COMPUTE_LINK_SIGN &	ab_link_sign =1

	return

__global__ void const RealA const const RealA const	QprodEven

__global__ void const RealA const const RealA const const RealA *const	QprodOdd

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB *const int int RealTypeId< RealA >::Type RealTypeId< RealA >::Type	accumu_coeff

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealTypeId< RealA >::Type RealA const	shortPEven

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB const int int RealTypeId< RealA >::Type RealTypeId< RealA >::Type RealA const RealA *const	shortPOdd

RealTypeId< RealA >::Type	mycoeff = CoeffSign<sig_positive,_oddBit ^ oddness_change>::result*coeff

RealA	COLOR_MAT_Z [ArrayLength< RealA >::result]

__global__ void const RealB const const RealA const	naikOprodEven

__global__ void const RealB const const RealA const const RealA *const	naikOprodOdd

__global__ void const RealB const const RealA const const RealA const int RealTypeId< RealA >::Type RealA const	outputEven

__global__ void const RealB const const RealA const const RealA const int RealTypeId< RealA >::Type RealA const RealA *const	outputOdd

RealA	de_link [ArrayLength< RealA >::result]

RealA	ef_link [ArrayLength< RealA >::result]

RealA	COLOR_MAT_U [ArrayLength< RealA >::result]

RealA	COLOR_MAT_V [ArrayLength< RealA >::result]

int	point_a

int	point_e

__global__ void const RealB const const RealA const	oprodEven

__global__ void const RealB const const RealA const const RealA const int RealA const	forceEven

__global__ void const RealB const const RealA const const RealA const int RealA const RealA *const	forceOdd

__global__ void const RealB const const RealA const const RealA const int RealA const RealA *const const int	threads

RealA	LINK_W [ArrayLength< RealA >::result]

COMPUTE_LINK_SIGN &	link_sign

int	bc_link_sign =1

int	de_link_sign =1

int	ef_link_sign =1

Typedef Documentation

typedef struct quda::fermion_force::hisq_kernel_param_s quda::fermion_force::hisq_kernel_param_t

Function Documentation

template<class Cmplx >

__device__ __host__ void quda::fermion_force::accumBothDerivatives	(	Matrix< Cmplx, 3 > *	result,
		const Matrix< Cmplx, 3 > &	left,
		const Matrix< Cmplx, 3 > &	right,
		const Matrix< Cmplx, 3 > &	outer_prod
	)

Definition at line 216 of file unitarize_force_quda.cu.

template<class Cmplx >

__device__ __host__ void quda::fermion_force::accumDerivatives	(	Matrix< Cmplx, 3 > *	result,
		const Matrix< Cmplx, 3 > &	left,
		const Matrix< Cmplx, 3 > &	right,
		const Matrix< Cmplx, 3 > &	outer_prod
	)

Definition at line 233 of file unitarize_force_quda.cu.

template<class T , class U >

__device__ void quda::fermion_force::addMatrixToField	(	const T *const	mat,
		int	dir,
		int	idx,
		U	coeff,
		T *const	field_even,
		T *const	field_odd,
		int	oddness
	)

inline

Definition at line 246 of file hisq_paths_force_quda.cu.

template<class T , class U >

__device__ void quda::fermion_force::addMatrixToField	(	const T *const	mat,
		int	idx,
		U	coeff,
		T *const	field_even,
		T *const	field_odd,
		int	oddness
	)

inline

Definition at line 266 of file hisq_paths_force_quda.cu.

quda::fermion_force::addMatrixToField	(	COLOR_MAT_W	,
		point_d	,
		accumu_coeff	,
		shortPEven	,
		shortPOdd	,
		1-	oddBit
	)

template<class T , class U >

__device__ void quda::fermion_force::addMatrixToField_test	(	const T *const	mat,
		int	idx,
		U	coeff,
		T *const	field_even,
		T *const	field_odd,
		int	oddness
	)

inline

Definition at line 285 of file hisq_paths_force_quda.cu.

quda::fermion_force::addMatrixToNewOprod	(	COLOR_MAT_W	,
		OPP_DIR(mu)	,
		new_sid	,
		mycoeff	,
		newOprodEven	,
		newOprodOdd	,
		oddBit
	)

quda::fermion_force::addMatrixToNewOprod	(	COLOR_MAT_W	,
		mu	,
		new_sid	,
		Sign< _oddBit^oddness_change >::result *	mycoeff,
		newOprodEven	,
		newOprodOdd	,
		oddBit
	)

quda::fermion_force::ADJ_MAT_MUL_ADJ_MAT	(	COLOR_MAT_X	,
		COLOR_MAT_Y	,
		COLOR_MAT_W
	)

template<class T >

__device__ void quda::fermion_force::adjointMatrix ( T * mat )

inline

Definition at line 134 of file hisq_paths_force_quda.cu.

quda::fermion_force::adjointMatrix ( ad_link )

template<class Real >

__device__ __host__ bool quda::fermion_force::checkAbsoluteError	(	Real	a,
		Real	b,
		Real	epsilon
	)

inline

Definition at line 259 of file unitarize_force_quda.cu.

template<class Real >

__device__ __host__ bool quda::fermion_force::checkRelativeError	(	Real	a,
		Real	b,
		Real	epsilon
	)

inline

Definition at line 268 of file unitarize_force_quda.cu.

template<class RealA , int oddBit>

__global__ void quda::fermion_force::do_one_link_term_kernel	(	const RealA *const	oprodEven,
		const RealA *const	oprodOdd,
		int	sig,
		typename RealTypeId< RealA >::Type	coeff,
		RealA *const	outputEven,
		RealA *const	outputOdd,
		const int	threads
	)

Definition at line 476 of file hisq_paths_force_quda.cu.

quda::fermion_force::FF_COMPUTE_NEW_FULL_IDX_MINUS_UPDATE	(	OPP_DIR(sig)	,
		new_mem_idx	,
		new_mem_idx
	)

quda::fermion_force::FF_COMPUTE_NEW_FULL_IDX_PLUS_UPDATE	(	OPP_DIR(mu)	,
		new_mem_idx	,
		new_mem_idx
	)

quda::fermion_force::FF_COMPUTE_NEW_FULL_IDX_PLUS_UPDATE	(	mymu	,
		new_mem_idx	,
		new_mem_idx
	)

quda::fermion_force::FF_COMPUTE_NEW_FULL_IDX_PLUS_UPDATE	(	mu	,
		new_mem_idx	,
		new_mem_idx
	)

template<class T >

__device__ __host__ T quda::fermion_force::getAbsMin	(	const T *const	array,
		int	size
	)

Definition at line 247 of file unitarize_force_quda.cu.

template<class Cmplx >

__global__ void quda::fermion_force::getUnitarizeForceField	(	const int	threads,
		const Cmplx *	link_even,
		const Cmplx *	link_odd,
		const Cmplx *	old_force_even,
		const Cmplx *	old_force_odd,
		Cmplx *	force_even,
		Cmplx *	force_odd,
		int *	unitarization_failed
	)

Definition at line 503 of file unitarize_force_quda.cu.

template<class Cmplx >

__device__ __host__ void quda::fermion_force::getUnitarizeForceSite	(	const Matrix< Cmplx, 3 > &	v,
		const Matrix< Cmplx, 3 > &	outer_prod,
		Matrix< Cmplx, 3 > *	result,
		int *	unitarization_failed
	)

Definition at line 441 of file unitarize_force_quda.cu.

template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>

__global__ void quda::fermion_force::HISQ_KERNEL_NAME	(	do_middle_link	,
		EXT
	)		const

template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>

__global__ void quda::fermion_force::HISQ_KERNEL_NAME	(	do_lepage_middle_link	,
		EXT
	)		const

template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>

__global__ void quda::fermion_force::HISQ_KERNEL_NAME	(	do_side_link	,
		EXT
	)		const

template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>

__global__ void quda::fermion_force::HISQ_KERNEL_NAME	(	do_side_link_short	,
		EXT
	)		const

template<class RealA , class RealB , SHORT sig_positive, SHORT mu_positive, SHORT _oddBit, int oddness_change>

__global__ void quda::fermion_force::HISQ_KERNEL_NAME	(	do_all_link	,
		EXT
	)		const

template<class RealA , class RealB , int oddBit>

__global__ void quda::fermion_force::HISQ_KERNEL_NAME	(	do_longlink	,
		EXT
	)		const

template<class RealA , class RealB , int oddBit>

__global__ void quda::fermion_force::HISQ_KERNEL_NAME	(	do_complete_force	,
		EXT
	)		const

quda::fermion_force::HISQ_LOAD_LINK	(	linkEven	,
		linkOdd	,
		mu	,
		new_sid	,
		ad_link	,
		oddBit
	)

quda::fermion_force::HISQ_LOAD_LINK	(	linkEven	,
		linkOdd	,
		mu	,
		point_b	,
		bc_link	,
		1-	oddBit
	)

quda::fermion_force::HISQ_LOAD_LINK	(	linkEven	,
		linkOdd	,
		sig	,
		new_sid	,
		LINK_W	,
		oddBit
	)

void quda::fermion_force::hisqCompleteForceCuda	(	const QudaGaugeParam &	param,
		const cudaGaugeField &	oprod,
		const cudaGaugeField &	link,
		cudaGaugeField *	force
	)

Definition at line 1917 of file hisq_paths_force_quda.cu.

void quda::fermion_force::hisqForceInitCuda ( QudaGaugeParam * param )

Definition at line 35 of file hisq_paths_force_quda.cu.

void quda::fermion_force::hisqLongLinkForceCuda	(	double	coeff,
		const QudaGaugeParam &	param,
		const cudaGaugeField &	oprod,
		const cudaGaugeField &	link,
		cudaGaugeField *	newOprod
	)

Definition at line 1943 of file hisq_paths_force_quda.cu.

void quda::fermion_force::hisqStaplesForceCuda	(	const double	path_coeff[6],
		const QudaGaugeParam &	param,
		const cudaGaugeField &	oprod,
		const cudaGaugeField &	link,
		cudaGaugeField *	newOprod
	)

Definition at line 1980 of file hisq_paths_force_quda.cu.

quda::fermion_force::if ( sid >=kparam. threads )

quda::fermion_force::if ( mu_positive )

Definition at line 174 of file hisq_paths_force_quda.cu.

quda::fermion_force::if ( sig_positive )

Definition at line 191 of file hisq_paths_force_quda.cu.

quda::fermion_force::if ( ! mu_positive )

Definition at line 225 of file hisq_paths_force_quda.cu.

quda::fermion_force::if ( QprevOdd = = NULL )

Definition at line 258 of file hisq_paths_force_quda.cu.

quda::fermion_force::if ( PmuOdd )

Definition at line 271 of file hisq_paths_force_quda.cu.

quda::fermion_force::if ( oddBit )

Definition at line 656 of file hisq_paths_force_quda.cu.

quda::fermion_force::if ( GOES_FORWARDS(sig) )

Definition at line 1086 of file hisq_paths_force_quda.cu.

quda::fermion_force::if ( sid >= threads )

template<int N, class T >

__device__ void quda::fermion_force::loadMatrixFromField	(	const T *const	field_even,
		const T *const	field_odd,
		int	dir,
		int	idx,
		T *const	mat,
		int	oddness,
		int	stride
	)

inline

Definition at line 159 of file hisq_paths_force_quda.cu.

template<class T >

__device__ void quda::fermion_force::loadMatrixFromField	(	const T *const	field_even,
		const T *const	field_odd,
		int	dir,
		int	idx,
		T *const	mat,
		int	oddness,
		int	stride
	)

inline

Definition at line 171 of file hisq_paths_force_quda.cu.

__device__ void quda::fermion_force::loadMatrixFromField	(	const float4 *const	field_even,
		const float4 *const	field_odd,
		int	dir,
		int	idx,
		float2 *const	mat,
		int	oddness,
		int	stride
	)

inline

Definition at line 181 of file hisq_paths_force_quda.cu.

template<class T >

__device__ void quda::fermion_force::loadMatrixFromField	(	const T *const	field_even,
		const T *const	field_odd,
		int	idx,
		T *const	mat,
		int	oddness,
		int	stride
	)

inline

Definition at line 200 of file hisq_paths_force_quda.cu.

quda::fermion_force::loadMatrixFromField	(	oprodEven	,
		oprodOdd	,
		point_c	,
		COLOR_MAT_Y	,
		oddBit	,
		hf.	color_matrix_stride
	)

quda::fermion_force::loadMatrixFromField	(	P3Even	,
		P3Odd	,
		new_sid	,
		COLOR_MAT_Y	,
		oddBit	,
		hf.	color_matrix_stride
	)

quda::fermion_force::loadMatrixFromField	(	QprodEven	,
		QprodOdd	,
		point_d	,
		COLOR_MAT_X	,
		1-	oddBit,
		hf.	color_matrix_stride
	)

quda::fermion_force::loadMatrixFromField	(	QprevEven	,
		QprevOdd	,
		point_d	,
		COLOR_MAT_X	,
		1-	oddBit,
		hf.	color_matrix_stride
	)

quda::fermion_force::loadMatrixFromField	(	oprodEven	,
		oprodOdd	,
		sig	,
		new_sid	,
		COLOR_MAT_X	,
		oddBit	,
		hf.	color_matrix_stride
	)

quda::fermion_force::MAT_MUL_MAT	(	bc_link	,
		COLOR_MAT_Y	,
		COLOR_MAT_Z
	)

quda::fermion_force::MAT_MUL_MAT	(	LINK_W	,
		COLOR_MAT_X	,
		COLOR_MAT_W
	)

quda::fermion_force::MATRIX_PRODUCT	(	bc_link	,
		COLOR_MAT_Y	,
		!	mu_positive,
		COLOR_MAT_W
	)

quda::fermion_force::MATRIX_PRODUCT	(	ab_link	,
		COLOR_MAT_W	,
		sig_positive	,
		COLOR_MAT_Y
	)

quda::fermion_force::MATRIX_PRODUCT	(	ad_link	,
		COLOR_MAT_Y	,
		mu_positive	,
		COLOR_MAT_W
	)

quda::fermion_force::MATRIX_PRODUCT	(	ab_link	,
		COLOR_MAT_Z	,
		sig_positive	,
		COLOR_MAT_Y
	)

quda::fermion_force::MATRIX_PRODUCT	(	ad_link	,
		COLOR_MAT_Y	,
		0	,
		COLOR_MAT_W
	)

__device__ float2 quda::fermion_force::operator*	(	float	a,
		const float2 &	b
	)

inline

Definition at line 79 of file hisq_paths_force_quda.cu.

__device__ double2 quda::fermion_force::operator*	(	double	a,
		const double2 &	b
	)

inline

Definition at line 84 of file hisq_paths_force_quda.cu.

__device__ const float2& quda::fermion_force::operator+=	(	float2 &	a,
		const float2 &	b
	)

inline

Definition at line 89 of file hisq_paths_force_quda.cu.

__device__ const double2& quda::fermion_force::operator+=	(	double2 &	a,
		const double2 &	b
	)

inline

Definition at line 96 of file hisq_paths_force_quda.cu.

__device__ const float4& quda::fermion_force::operator+=	(	float4 &	a,
		const float4 &	b
	)

inline

Definition at line 103 of file hisq_paths_force_quda.cu.

template<class Cmplx >

__device__ __host__ void quda::fermion_force::reciprocalRoot	(	Matrix< Cmplx, 3 > *	res,
		DerivativeCoefficients< typename RealTypeId< Cmplx >::Type > *	deriv_coeffs,
		typename RealTypeId< Cmplx >::Type	f[3],
		Matrix< Cmplx, 3 > &	q,
		int *	unitarization_failed
	)

Definition at line 281 of file unitarize_force_quda.cu.

quda::fermion_force::RECONSTRUCT_SITE_LINK	(	ad_link	,
		ad_link_sign
	)

quda::fermion_force::RECONSTRUCT_SITE_LINK	(	bc_link	,
		bc_link_sign
	)

quda::fermion_force::RECONSTRUCT_SITE_LINK	(	LINK_W	,
		link_sign
	)

template<typename T >

__device__ void quda::fermion_force::reconstructSign	(	int *const	sign,
		int	dir,
		const T	i[4]
	)

Definition at line 431 of file hisq_paths_force_quda.cu.

void quda::fermion_force::setUnitarizeForceConstants	(	double	unitarize_eps,
		double	hisq_force_filter,
		double	max_det_error,
		bool	allow_svd,
		bool	svd_only,
		double	svd_rel_error,
		double	svd_abs_error
	)

Definition at line 41 of file unitarize_force_quda.cu.

quda::fermion_force::storeMatrixToField	(	COLOR_MAT_Y	,
		new_sid	,
		P3Even	,
		P3Odd	,
		oddBit
	)

template<class T >

__device__ void quda::fermion_force::storeMatrixToField	(	const T *const	mat,
		int	dir,
		int	idx,
		T *const	field_even,
		T *const	field_odd,
		int	oddness
	)

inline

Definition at line 308 of file hisq_paths_force_quda.cu.

template<class T >

__device__ void quda::fermion_force::storeMatrixToField	(	const T *const	mat,
		int	idx,
		T *const	field_even,
		T *const	field_odd,
		int	oddness
	)

inline

Definition at line 327 of file hisq_paths_force_quda.cu.

template<class T , class U >

__device__ void quda::fermion_force::storeMatrixToMomentumField	(	const T *const	mat,
		int	dir,
		int	idx,
		U	coeff,
		T *const	mom_even,
		T *const	mom_odd,
		int	oddness
	)

inline

Definition at line 346 of file hisq_paths_force_quda.cu.

quda::fermion_force::storeMatrixToMomentumField	(	COLOR_MAT_W	,
		sig	,
		sid	,
		coeff	,
		forceEven	,
		forceOdd	,
		oddBit
	)

void quda::fermion_force::unitarizeForceCPU	(	const QudaGaugeParam &	param,
		cpuGaugeField &	cpuOldForce,
		cpuGaugeField &	cpuGauge,
		cpuGaugeField *	cpuNewForce
	)

Definition at line 544 of file unitarize_force_quda.cu.

void quda::fermion_force::unitarizeForceCuda	(	const QudaGaugeParam &	param,
		cudaGaugeField &	cudaOldForce,
		cudaGaugeField &	cudaGauge,
		cudaGaugeField *	cudaNewForce,
		int *	unitarization_failed
	)

Definition at line 685 of file unitarize_force_quda.cu.

Variable Documentation

RealA quda::fermion_force::ab_link

Definition at line 131 of file hisq_paths_force_quda.cu.

int quda::fermion_force::ab_link_nbr_idx = (sig_positive) ? new_sid : point_b

Definition at line 148 of file hisq_paths_force_quda.cu.

int quda::fermion_force::ab_link_sign =1

Definition at line 235 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const SHORT SHORT RealTypeId< RealA >::Type RealTypeId< RealA >::Type quda::fermion_force::accumu_coeff

Definition at line 539 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::ad_link

Definition at line 133 of file hisq_paths_force_quda.cu.

int quda::fermion_force::ad_link_nbr_idx

Definition at line 148 of file hisq_paths_force_quda.cu.

COMPUTE_LINK_SIGN & quda::fermion_force::ad_link_sign

Definition at line 187 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::bc_link

Definition at line 132 of file hisq_paths_force_quda.cu.

int quda::fermion_force::bc_link_nbr_idx

Definition at line 148 of file hisq_paths_force_quda.cu.

int quda::fermion_force::bc_link_sign =1

Definition at line 128 of file hisq_paths_force_quda.cu.

RealTypeId< RealA >::Type quda::fermion_force::coeff = (oddBit==1) ? -1 : 1

Definition at line 97 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::COLOR_MAT_U

Definition at line 1062 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::COLOR_MAT_V

Definition at line 1063 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::COLOR_MAT_W

Definition at line 135 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::COLOR_MAT_X

Definition at line 137 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::COLOR_MAT_Y

Definition at line 136 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::COLOR_MAT_Z

Definition at line 830 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::de_link

Definition at line 1052 of file hisq_paths_force_quda.cu.

int quda::fermion_force::de_link_sign =1

Definition at line 1058 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::ef_link

Definition at line 1053 of file hisq_paths_force_quda.cu.

int quda::fermion_force::ef_link_sign =1

Definition at line 1059 of file hisq_paths_force_quda.cu.

quda::fermion_force::else

Initial value:

{

mymu = OPP_DIR(mu)

Definition at line 177 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const int RealA *const quda::fermion_force::forceEven

Definition at line 1147 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const int RealA *const RealA *const quda::fermion_force::forceOdd

Definition at line 1147 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const int RealTypeId< RealA >::Type RealA *const RealA *const hisq_kernel_param_t quda::fermion_force::kparam

Initial value:

{
  int oddBit = _oddBit

Definition at line 107 of file hisq_paths_force_quda.cu.

COMPUTE_LINK_SIGN & quda::fermion_force::link_sign

Definition at line 1187 of file hisq_paths_force_quda.cu.

RealA quda::fermion_force::LINK_W

Definition at line 1181 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const quda::fermion_force::linkEven

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const quda::fermion_force::linkOdd

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const SHORT SHORT quda::fermion_force::mu

Definition at line 97 of file hisq_paths_force_quda.cu.

const RealTypeId< RealA >::Type & quda::fermion_force::mycoeff = CoeffSign<sig_positive,_oddBit ^ oddness_change>::result*coeff

Definition at line 615 of file hisq_paths_force_quda.cu.

int quda::fermion_force::mymu

Definition at line 149 of file hisq_paths_force_quda.cu.

int quda::fermion_force::mysig

Definition at line 190 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const quda::fermion_force::naikOprodEven

Definition at line 1008 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const quda::fermion_force::naikOprodOdd

Definition at line 1008 of file hisq_paths_force_quda.cu.

int quda::fermion_force::new_mem_idx = X

Definition at line 124 of file hisq_paths_force_quda.cu.

int quda::fermion_force::new_sid = sid

Definition at line 171 of file hisq_paths_force_quda.cu.

int quda::fermion_force::new_x = x[0]

Definition at line 123 of file hisq_paths_force_quda.cu.

texture<int4, 1> quda::fermion_force::newOprod0TexDouble

Definition at line 29 of file hisq_paths_force_quda.cu.

texture<float2, 1, cudaReadModeElementType> quda::fermion_force::newOprod0TexSingle

Definition at line 31 of file hisq_paths_force_quda.cu.

texture<int4, 1> quda::fermion_force::newOprod1TexDouble

Definition at line 30 of file hisq_paths_force_quda.cu.

texture<float2, 1, cudaReadModeElementType> quda::fermion_force::newOprod1TexSingle

Definition at line 32 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const SHORT SHORT RealTypeId< RealA >::Type RealTypeId< RealA >::Type RealA *const RealA *const RealA *const quda::fermion_force::newOprodEven

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const SHORT SHORT RealTypeId< RealA >::Type RealTypeId< RealA >::Type RealA *const RealA *const RealA *const RealA *const quda::fermion_force::newOprodOdd

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const quda::fermion_force::oprodEven

Definition at line 1147 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const quda::fermion_force::oprodOdd

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const int RealTypeId< RealA >::Type RealA *const quda::fermion_force::outputEven

Definition at line 1008 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const int RealTypeId< RealA >::Type RealA *const RealA *const quda::fermion_force::outputOdd

Definition at line 1008 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int RealTypeId< RealA >::Type RealA *const quda::fermion_force::P3Even

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const quda::fermion_force::P3Odd

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int RealTypeId< RealA >::Type RealA *const quda::fermion_force::PmuEven

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int RealTypeId< RealA >::Type RealA *const RealA *const quda::fermion_force::PmuOdd

Definition at line 97 of file hisq_paths_force_quda.cu.

int quda::fermion_force::point_a

Definition at line 1071 of file hisq_paths_force_quda.cu.

int quda::fermion_force::point_b = (new_mem_idx >> 1)

Definition at line 147 of file hisq_paths_force_quda.cu.

const int & quda::fermion_force::point_c = (new_mem_idx >> 1)

Definition at line 147 of file hisq_paths_force_quda.cu.

int quda::fermion_force::point_d = (new_mem_idx >> 1)

Definition at line 147 of file hisq_paths_force_quda.cu.

int quda::fermion_force::point_e

Definition at line 1071 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int RealTypeId< RealA >::Type RealA *const RealA *const RealA *const RealA *const RealA *const quda::fermion_force::QmuEven

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int RealTypeId< RealA >::Type RealA *const RealA *const RealA *const RealA *const RealA *const RealA *const quda::fermion_force::QmuOdd

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const quda::fermion_force::QprevEven

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const quda::fermion_force::QprevOdd

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const quda::fermion_force::QprodEven

Definition at line 539 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const quda::fermion_force::QprodOdd

Definition at line 539 of file hisq_paths_force_quda.cu.

quda::fermion_force::return

Definition at line 501 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const SHORT SHORT RealTypeId< RealA >::Type RealTypeId< RealA >::Type RealA *const quda::fermion_force::shortPEven

Definition at line 539 of file hisq_paths_force_quda.cu.

__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const SHORT SHORT RealTypeId< RealA >::Type RealTypeId< RealA >::Type RealA *const RealA *const quda::fermion_force::shortPOdd

Definition at line 539 of file hisq_paths_force_quda.cu.

int quda::fermion_force::sid = blockIdx.x * blockDim.x + threadIdx.x

Definition at line 111 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const int quda::fermion_force::sig

Definition at line 97 of file hisq_paths_force_quda.cu.

__global__ void const RealB *const const RealA *const const RealA *const int RealA *const RealA *const const int quda::fermion_force::threads

Initial value:

{
  int sid = blockIdx.x * blockDim.x + threadIdx.x

Definition at line 1152 of file hisq_paths_force_quda.cu.

int quda::fermion_force::x = z1 - z2*D2

Definition at line 113 of file hisq_paths_force_quda.cu.

int quda::fermion_force::X = 2*sid + x1odd

Definition at line 165 of file hisq_paths_force_quda.cu.

int quda::fermion_force::x1h = sid - z1*D1h

Definition at line 115 of file hisq_paths_force_quda.cu.

int quda::fermion_force::x1odd = (x[1] + x[2] + x[3] + oddBit) & 1

Definition at line 120 of file hisq_paths_force_quda.cu.

int quda::fermion_force::z1 = sid/D1h

Definition at line 114 of file hisq_paths_force_quda.cu.

int quda::fermion_force::z2 = z1/D2

Definition at line 116 of file hisq_paths_force_quda.cu.

Classes

Typedefs

Functions

Variables

Typedef Documentation

Function Documentation

Variable Documentation