#include <read_gauge.h>
#include <gauge_field.h>
#include <hisq_force_quda.h>
#include <hw_quda.h>
#include <hisq_force_macros.h>
#include <utility>
#include "hisq_paths_force_core.h"

Go to the source code of this file.

Classes
struct	quda::fermion_force::hisq_kernel_param_s

struct	quda::fermion_force::PathCoefficients< Real >

struct	RealTypeId< T >

struct	quda::fermion_force::RealTypeId< float2 >

struct	quda::fermion_force::RealTypeId< double2 >

struct	quda::fermion_force::CoeffSign< pos_dir, odd_lattice >

struct	quda::fermion_force::CoeffSign< 0, 1 >

struct	quda::fermion_force::CoeffSign< 0, 0 >

struct	quda::fermion_force::CoeffSign< 1, 1 >

struct	quda::fermion_force::Sign< odd_lattice >

struct	quda::fermion_force::Sign< 1 >

struct	quda::fermion_force::ArrayLength< RealX >

struct	quda::fermion_force::ArrayLength< float4 >

class	quda::fermion_force::MiddleLink< RealA, RealB >

class	quda::fermion_force::LepageMiddleLink< RealA, RealB >

class	quda::fermion_force::SideLink< RealA, RealB >

class	quda::fermion_force::SideLinkShort< RealA, RealB >

class	quda::fermion_force::AllLink< RealA, RealB >

class	quda::fermion_force::OneLinkTerm< RealA, RealB >

class	quda::fermion_force::LongLinkTerm< RealA, RealB >

class	quda::fermion_force::CompleteForce< RealA, RealB >

Namespaces
namespace	quda

namespace	quda::fermion_force

Macros
#define	COMPILE_HISQ_DP_18

#define	COMPILE_HISQ_DP_12

#define	COMPILE_HISQ_SP_18

#define	COMPILE_HISQ_SP_12

#define	HISQ_SITE_MATRIX_LOAD_TEX 1

#define	HISQ_NEW_OPROD_LOAD_TEX 1

#define	CONJ_INDEX(i, j) j*3 + i

#define	addMatrixToNewOprod(mat,dir, idx, coeff, field_even, field_odd, oddness)

#define	DD_CONCAT(n, r) n ## r ## kernel

#define	HISQ_KERNEL_NAME(a, b) DD_CONCAT(a,b)

#define	NEWOPROD_EVEN_TEX newOprod0TexDouble

#define	NEWOPROD_ODD_TEX newOprod1TexDouble

#define	LOAD_TEX_ENTRY(tex, field, idx) READ_DOUBLE2_TEXTURE(tex, field, idx)

#define	PRECISION 0

#define	RECON 18

#define	HISQ_LOAD_LINK(linkEven, linkOdd, dir, idx, var, oddness) HISQ_LOAD_MATRIX_18_DOUBLE_TEX((oddness)?siteLink1TexDouble:siteLink0TexDouble, (oddness)?linkOdd:linkEven, dir, idx, var, hf.site_ga_stride)

#define	COMPUTE_LINK_SIGN(sign, dir, x)

#define	RECONSTRUCT_SITE_LINK(var, sign)

#define	PRECISION 0

#define	RECON 12

#define	HISQ_LOAD_LINK(linkEven, linkOdd, dir, idx, var, oddness) HISQ_LOAD_MATRIX_12_DOUBLE_TEX((oddness)?siteLink1TexDouble:siteLink0TexDouble, (oddness)?linkOdd:linkEven,dir, idx, var, hf.site_ga_stride)

#define	COMPUTE_LINK_SIGN(sign, dir, x) reconstructSign(sign, dir, x)

#define	RECONSTRUCT_SITE_LINK(var, sign) FF_RECONSTRUCT_LINK_12(var, sign)

#define	NEWOPROD_EVEN_TEX newOprod0TexSingle

#define	NEWOPROD_ODD_TEX newOprod1TexSingle

#define	LOAD_TEX_ENTRY(tex, field, idx) tex1Dfetch(tex,idx)

#define	PRECISION 1

#define	RECON 18

#define	HISQ_LOAD_LINK(linkEven, linkOdd, dir, idx, var, oddness) HISQ_LOAD_MATRIX_18_SINGLE_TEX((oddness)?siteLink1TexSingle:siteLink0TexSingle, dir, idx, var, hf.site_ga_stride)

#define	COMPUTE_LINK_SIGN(sign, dir, x)

#define	RECONSTRUCT_SITE_LINK(var, sign)

#define	PRECISION 1

#define	RECON 12

#define	HISQ_LOAD_LINK(linkEven, linkOdd, dir, idx, var, oddness) HISQ_LOAD_MATRIX_12_SINGLE_TEX((oddness)?siteLink1TexSingle_recon:siteLink0TexSingle_recon, dir, idx, var, hf.site_ga_stride)

#define	COMPUTE_LINK_SIGN(sign, dir, x) reconstructSign(sign, dir, x)

#define	RECONSTRUCT_SITE_LINK(var, sign) FF_RECONSTRUCT_LINK_12(var, sign)

#define	CALL_ARGUMENTS(typeA, typeB)

#define	CALL_MIDDLE_LINK_KERNEL(sig_sign, mu_sign)

#define	CALL_ARGUMENTS(typeA, typeB)

#define	CALL_MIDDLE_LINK_KERNEL(sig_sign, mu_sign)

#define	CALL_ARGUMENTS(typeA, typeB)

#define	CALL_SIDE_LINK_KERNEL(sig_sign, mu_sign)

#define	CALL_ARGUMENTS(typeA, typeB)

#define	CALL_SIDE_LINK_KERNEL(sig_sign, mu_sign)

#define	CALL_ARGUMENTS(typeA, typeB)

#define	CALL_ALL_LINK_KERNEL(sig_sign, mu_sign)

#define	CALL_ARGUMENTS(typeA, typeB)

#define	CALL_ARGUMENTS(typeA, typeB)

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

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

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

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

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

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

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

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

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)

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)

__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)

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)

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)

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)

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)

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)

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)

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)

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

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)

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 RealA *const oprodEven

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

	quda::fermion_force::if (mu_positive)

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

	quda::fermion_force::if (sig_positive)

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

	quda::fermion_force::if (!mu_positive)

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

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

	quda::fermion_force::if (PmuOdd)

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

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

	quda::fermion_force::adjointMatrix (ad_link)

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 RealA *const oprodEven

	quda::fermion_force::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	quda::fermion_force::HISQ_KERNEL_NAME (do_side_link, EXT)(const RealA *const P3Even

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

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

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

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

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

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

	quda::fermion_force::if (oddBit)

	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::HISQ_KERNEL_NAME (do_all_link, EXT)(const RealA *const oprodEven

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

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

	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::RECONSTRUCT_SITE_LINK (bc_link, bc_link_sign)

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

__global__ void const RealA const const RealA const const RealA const const RealB const const RealB *const int int RealTypeId< RealA >::Type	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::kparam

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

int	quda::fermion_force::x [4] = z1 - z2*D2

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

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

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

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

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

int	quda::fermion_force::new_mem_idx = X

RealA	quda::fermion_force::ab_link [ArrayLength< RealA >::result]

RealA	quda::fermion_force::bc_link [ArrayLength< RealA >::result]

RealA	quda::fermion_force::ad_link [ArrayLength< RealA >::result]

RealA	quda::fermion_force::COLOR_MAT_W [ArrayLength< RealA >::result]

RealA	quda::fermion_force::COLOR_MAT_Y [ArrayLength< RealA >::result]

RealA	quda::fermion_force::COLOR_MAT_X [ArrayLength< RealA >::result]

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

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

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

int	quda::fermion_force::ad_link_nbr_idx

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

int	quda::fermion_force::bc_link_nbr_idx

int	quda::fermion_force::mymu

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

int	quda::fermion_force::new_sid = sid

	quda::fermion_force::else

COMPUTE_LINK_SIGN &	quda::fermion_force::ad_link_sign

int	quda::fermion_force::mysig

COMPUTE_LINK_SIGN &	quda::fermion_force::ab_link_sign =1

	quda::fermion_force::return

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

__global__ void const RealA const const RealA const const RealA *const	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::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	quda::fermion_force::shortPOdd

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

RealA	quda::fermion_force::COLOR_MAT_Z [ArrayLength< RealA >::result]

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

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

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

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

RealA	quda::fermion_force::de_link [ArrayLength< RealA >::result]

RealA	quda::fermion_force::ef_link [ArrayLength< RealA >::result]

RealA	quda::fermion_force::COLOR_MAT_U [ArrayLength< RealA >::result]

RealA	quda::fermion_force::COLOR_MAT_V [ArrayLength< RealA >::result]

int	quda::fermion_force::point_a

int	quda::fermion_force::point_e

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

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

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

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

RealA	quda::fermion_force::LINK_W [ArrayLength< RealA >::result]

COMPUTE_LINK_SIGN &	quda::fermion_force::link_sign

int	quda::fermion_force::bc_link_sign =1

int	quda::fermion_force::de_link_sign =1

int	quda::fermion_force::ef_link_sign =1

Macro Definition Documentation

#define addMatrixToNewOprod	(	mat,
		dir,
		idx,
		coeff,
		field_even,
		field_odd,
		oddness
	)

Value:

do { \
      RealA* const field = (oddness)?field_odd: field_even;             \
      RealA value[9];                                                   \
      value[0] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9); \
      value[1] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9 + hf.color_matrix_stride); \
      value[2] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9 + 2*hf.color_matrix_stride); \
      value[3] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9 + 3*hf.color_matrix_stride); \
      value[4] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9 + 4*hf.color_matrix_stride); \
      value[5] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9 + 5*hf.color_matrix_stride); \
      value[6] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9 + 6*hf.color_matrix_stride); \
      value[7] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9 + 7*hf.color_matrix_stride); \
      value[8] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*hf.color_matrix_stride*9 + 8*hf.color_matrix_stride); \
      field[idx + dir*hf.color_matrix_stride*9]          = value[0] + coeff*mat[0]; \
      field[idx + dir*hf.color_matrix_stride*9 + hf.color_matrix_stride]     = value[1] + coeff*mat[1]; \
      field[idx + dir*hf.color_matrix_stride*9 + hf.color_matrix_stride*2]   = value[2] + coeff*mat[2]; \
      field[idx + dir*hf.color_matrix_stride*9 + hf.color_matrix_stride*3]   = value[3] + coeff*mat[3]; \
      field[idx + dir*hf.color_matrix_stride*9 + hf.color_matrix_stride*4]   = value[4] + coeff*mat[4]; \
      field[idx + dir*hf.color_matrix_stride*9 + hf.color_matrix_stride*5]   = value[5] + coeff*mat[5]; \
      field[idx + dir*hf.color_matrix_stride*9 + hf.color_matrix_stride*6]   = value[6] + coeff*mat[6]; \
      field[idx + dir*hf.color_matrix_stride*9 + hf.color_matrix_stride*7]   = value[7] + coeff*mat[7]; \
      field[idx + dir*hf.color_matrix_stride*9 + hf.color_matrix_stride*8]   = value[8] + coeff*mat[8]; \
    }while(0)

Definition at line 217 of file hisq_paths_force_quda.cu.

#define CALL_ALL_LINK_KERNEL	(	sig_sign,
		mu_sign
	)

Definition at line 1294 of file hisq_paths_force_quda.cu.

#define CALL_ARGUMENTS	(	typeA,
		typeB
	)

Value:

<<<tp.grid, tp.block>>>         \
      ((typeA*)oprod.Even_p(), (typeA*)oprod.Odd_p(),                   \
       (typeA*)Qprev_even, (typeA*)Qprev_odd,                           \
       (typeB*)link.Even_p(), (typeB*)link.Odd_p(),                     \
       sig, mu, coeff,                                                  \
       (typeA*)Pmu.Even_p(), (typeA*)Pmu.Odd_p(),                       \
       (typeA*)P3.Even_p(), (typeA*)P3.Odd_p(),                         \
       (typeA*)Qmu.Even_p(), (typeA*)Qmu.Odd_p(),                       \
       (typeA*)newOprod.Even_p(), (typeA*)newOprod.Odd_p(), kparam)

Definition at line 1624 of file hisq_paths_force_quda.cu.

#define CALL_ARGUMENTS	(	typeA,
		typeB
	)

Value:

<<<tp.grid, tp.block>>>         \
      ((typeA*)oprod.Even_p(), (typeA*)oprod.Odd_p(),                   \
       (typeA*)Qprev.Even_p(), (typeA*)Qprev.Odd_p(),                   \
       (typeB*)link.Even_p(), (typeB*)link.Odd_p(),                     \
       sig, mu, coeff,                                                  \
       (typeA*)P3.Even_p(), (typeA*)P3.Odd_p(),                         \
       (typeA*)newOprod.Even_p(), (typeA*)newOprod.Odd_p(),             \
       kparam)

Definition at line 1624 of file hisq_paths_force_quda.cu.

#define CALL_ARGUMENTS	(	typeA,
		typeB
	)

Value:

<<<tp.grid, tp.block>>>         \
      ((typeA*)P3.Even_p(), (typeA*)P3.Odd_p(),                         \
       (typeA*)oprod.Even_p(),  (typeA*)oprod.Odd_p(),                  \
       (typeB*)link.Even_p(), (typeB*)link.Odd_p(),                     \
       sig, mu,                                                         \
       coeff,                   \
       (typename RealTypeId<typeA>::Type) accumu_coeff,                 \
       (typeA*)shortP.Even_p(), (typeA*)shortP.Odd_p(),                 \
       (typeA*)newOprod.Even_p(), (typeA*)newOprod.Odd_p(),             \
       kparam)

Definition at line 1624 of file hisq_paths_force_quda.cu.

#define CALL_ARGUMENTS	(	typeA,
		typeB
	)

Value:

<<<tp.grid, tp.block>>>         \
      ((typeA*)P3.Even_p(), (typeA*)P3.Odd_p(),                         \
       (typeB*)link.Even_p(), (typeB*)link.Odd_p(),                     \
       sig, mu, (typename RealTypeId<typeA>::Type) coeff,               \
       (typeA*)newOprod.Even_p(), (typeA*)newOprod.Odd_p(), kparam)

Definition at line 1624 of file hisq_paths_force_quda.cu.

#define CALL_ARGUMENTS	(	typeA,
		typeB
	)

Value:

<<<tp.grid, tp.block>>>         \
      ((typeA*)oprod.Even_p(), (typeA*)oprod.Odd_p(),                   \
       (typeA*)Qprev.Even_p(), (typeA*)Qprev.Odd_p(),                   \
       (typeB*)link.Even_p(), (typeB*)link.Odd_p(), sig,  mu,           \
       (typename RealTypeId<typeA>::Type)coeff,                         \
       (typename RealTypeId<typeA>::Type)accumu_coeff,                  \
       (typeA*)shortP.Even_p(),(typeA*)shortP.Odd_p(),                  \
       (typeA*)newOprod.Even_p(), (typeA*)newOprod.Odd_p(), kparam)

Definition at line 1624 of file hisq_paths_force_quda.cu.

#define CALL_ARGUMENTS	(	typeA,
		typeB
	)

Value:

<<<tp.grid,tp.block>>>          \
        ((typeB*)link.Even_p(), (typeB*)link.Odd_p(),                   \
       (typeA*)naikOprod.Even_p(),  (typeA*)naikOprod.Odd_p(),          \
       sig, naik_coeff,                                                 \
       (typeA*)output.Even_p(), (typeA*)output.Odd_p(),                 \
       kparam);

Definition at line 1624 of file hisq_paths_force_quda.cu.

#define CALL_ARGUMENTS	(	typeA,
		typeB
	)

Value:

<<<tp.grid, tp.block>>>         \
      ((typeB*)link.Even_p(), (typeB*)link.Odd_p(),                     \
       (typeA*)oprod.Even_p(), (typeA*)oprod.Odd_p(),                   \
       sig,                                                             \
       (typeA*)mom.Even_p(), (typeA*)mom.Odd_p(),                       \
       X[0] * X[1] * X[2] * X[3]/2);

Definition at line 1624 of file hisq_paths_force_quda.cu.

#define CALL_MIDDLE_LINK_KERNEL	(	sig_sign,
		mu_sign
	)

Definition at line 862 of file hisq_paths_force_quda.cu.

#define CALL_MIDDLE_LINK_KERNEL	(	sig_sign,
		mu_sign
	)

Definition at line 862 of file hisq_paths_force_quda.cu.

#define CALL_SIDE_LINK_KERNEL	(	sig_sign,
		mu_sign
	)

Definition at line 1148 of file hisq_paths_force_quda.cu.

#define CALL_SIDE_LINK_KERNEL	(	sig_sign,
		mu_sign
	)

Definition at line 1148 of file hisq_paths_force_quda.cu.

#define COMPILE_HISQ_DP_12

Definition at line 11 of file hisq_paths_force_quda.cu.

#define COMPILE_HISQ_DP_18

Definition at line 10 of file hisq_paths_force_quda.cu.

#define COMPILE_HISQ_SP_12

Definition at line 13 of file hisq_paths_force_quda.cu.

#define COMPILE_HISQ_SP_18

Definition at line 12 of file hisq_paths_force_quda.cu.

#define COMPUTE_LINK_SIGN	(	sign,
		dir,
		x
	)

Definition at line 592 of file hisq_paths_force_quda.cu.

#define COMPUTE_LINK_SIGN	(	sign,
		dir,
		x
	)	reconstructSign(sign, dir, x)

Definition at line 592 of file hisq_paths_force_quda.cu.

#define COMPUTE_LINK_SIGN	(	sign,
		dir,
		x
	)

Definition at line 592 of file hisq_paths_force_quda.cu.

#define COMPUTE_LINK_SIGN	(	sign,
		dir,
		x
	)	reconstructSign(sign, dir, x)

Definition at line 592 of file hisq_paths_force_quda.cu.

#define CONJ_INDEX	(	i,
		j
	)	j*3 + i

#define DD_CONCAT	(	n,
		r
	)	n ## r ## kernel

Definition at line 507 of file hisq_paths_force_quda.cu.

#define HISQ_KERNEL_NAME	(	a,
		b
	)	DD_CONCAT(a,b)

Definition at line 509 of file hisq_paths_force_quda.cu.

#define HISQ_LOAD_LINK	(	linkEven,
		linkOdd,
		dir,
		idx,
		var,
		oddness
	)	HISQ_LOAD_MATRIX_18_DOUBLE_TEX((oddness)?siteLink1TexDouble:siteLink0TexDouble, (oddness)?linkOdd:linkEven, dir, idx, var, hf.site_ga_stride)

Definition at line 588 of file hisq_paths_force_quda.cu.

#define HISQ_LOAD_LINK	(	linkEven,
		linkOdd,
		dir,
		idx,
		var,
		oddness
	)	HISQ_LOAD_MATRIX_12_DOUBLE_TEX((oddness)?siteLink1TexDouble:siteLink0TexDouble, (oddness)?linkOdd:linkEven,dir, idx, var, hf.site_ga_stride)

Definition at line 588 of file hisq_paths_force_quda.cu.

#define HISQ_LOAD_LINK	(	linkEven,
		linkOdd,
		dir,
		idx,
		var,
		oddness
	)	HISQ_LOAD_MATRIX_18_SINGLE_TEX((oddness)?siteLink1TexSingle:siteLink0TexSingle, dir, idx, var, hf.site_ga_stride)

Definition at line 588 of file hisq_paths_force_quda.cu.

#define HISQ_LOAD_LINK	(	linkEven,
		linkOdd,
		dir,
		idx,
		var,
		oddness
	)	HISQ_LOAD_MATRIX_12_SINGLE_TEX((oddness)?siteLink1TexSingle_recon:siteLink0TexSingle_recon, dir, idx, var, hf.site_ga_stride)

Definition at line 588 of file hisq_paths_force_quda.cu.

#define HISQ_NEW_OPROD_LOAD_TEX 1

Definition at line 17 of file hisq_paths_force_quda.cu.

#define HISQ_SITE_MATRIX_LOAD_TEX 1

Definition at line 16 of file hisq_paths_force_quda.cu.

#define LOAD_TEX_ENTRY	(	tex,
		field,
		idx
	)	READ_DOUBLE2_TEXTURE(tex, field, idx)

Definition at line 562 of file hisq_paths_force_quda.cu.

#define LOAD_TEX_ENTRY	(	tex,
		field,
		idx
	)	tex1Dfetch(tex,idx)

Definition at line 562 of file hisq_paths_force_quda.cu.

#define NEWOPROD_EVEN_TEX newOprod0TexDouble

Definition at line 558 of file hisq_paths_force_quda.cu.

#define NEWOPROD_EVEN_TEX newOprod0TexSingle

Definition at line 558 of file hisq_paths_force_quda.cu.

#define NEWOPROD_ODD_TEX newOprod1TexDouble

Definition at line 559 of file hisq_paths_force_quda.cu.

#define NEWOPROD_ODD_TEX newOprod1TexSingle

Definition at line 559 of file hisq_paths_force_quda.cu.

#define PRECISION 0

Definition at line 585 of file hisq_paths_force_quda.cu.

#define PRECISION 0

Definition at line 585 of file hisq_paths_force_quda.cu.

#define PRECISION 1

Definition at line 585 of file hisq_paths_force_quda.cu.

#define PRECISION 1

Definition at line 585 of file hisq_paths_force_quda.cu.

#define RECON 18

Definition at line 586 of file hisq_paths_force_quda.cu.

#define RECON 12

Definition at line 586 of file hisq_paths_force_quda.cu.

#define RECON 18

Definition at line 586 of file hisq_paths_force_quda.cu.

#define RECON 12

Definition at line 586 of file hisq_paths_force_quda.cu.

#define RECONSTRUCT_SITE_LINK	(	var,
		sign
	)

Definition at line 593 of file hisq_paths_force_quda.cu.

#define RECONSTRUCT_SITE_LINK	(	var,
		sign
	)	FF_RECONSTRUCT_LINK_12(var, sign)

Definition at line 593 of file hisq_paths_force_quda.cu.

#define RECONSTRUCT_SITE_LINK	(	var,
		sign
	)

Definition at line 593 of file hisq_paths_force_quda.cu.

#define RECONSTRUCT_SITE_LINK	(	var,
		sign
	)	FF_RECONSTRUCT_LINK_12(var, sign)

Definition at line 593 of file hisq_paths_force_quda.cu.

Classes

Namespaces

Macros

Typedefs

Functions

Variables

Macro Definition Documentation