QUDA  v0.5.0
A library for QCD on GPUs
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
Namespaces | Macros | Functions | Variables
llfat_quda.cu File Reference
#include <stdio.h>
#include <cuda_runtime.h>
#include <cuda.h>
#include <quda_internal.h>
#include <llfat_quda.h>
#include <read_gauge.h>
#include "gauge_field.h"
#include <force_common.h>
#include "llfat_core.h"

Go to the source code of this file.

Namespaces

namespace  quda
 

Macros

#define SITE_MATRIX_LOAD_TEX   0
 
#define MULINK_LOAD_TEX   1
 
#define FATLINK_LOAD_TEX   1
 
#define BLOCK_DIM   64
 
#define WRITE_FAT_MATRIX(gauge, dir, idx)
 
#define WRITE_STAPLE_MATRIX(gauge, idx)
 
#define SCALAR_MULT_SU3_MATRIX(a, b, c)
 
#define LOAD_MATRIX_12_SINGLE_DECLARE(gauge, dir, idx, var, stride)
 
#define LOAD_MATRIX_12_SINGLE_TEX_DECLARE(gauge, dir, idx, var, stride)
 
#define LOAD_MATRIX_18_SINGLE_DECLARE(gauge, dir, idx, var, stride)
 
#define LOAD_MATRIX_18_SINGLE_TEX_DECLARE(gauge, dir, idx, var, stride)
 
#define LOAD_MATRIX_18_DOUBLE_DECLARE(gauge, dir, idx, var, stride)
 
#define LOAD_MATRIX_18_DOUBLE_TEX_DECLARE(gauge_tex, gauge, dir, idx, var, stride)
 
#define LOAD_MATRIX_12_DOUBLE_DECLARE(gauge, dir, idx, var, stride)
 
#define LOAD_MATRIX_12_DOUBLE_TEX_DECLARE(gauge_tex, gauge, dir, idx, var, stride)
 
#define LLFAT_ADD_SU3_MATRIX(ma, mb, mc)
 
#define LLFAT_CONCAT(a, b)   a##b##Kernel
 
#define LLFAT_CONCAT_EX(a, b)   a##b##Kernel_ex
 
#define LLFAT_KERNEL(a, b)   LLFAT_CONCAT(a,b)
 
#define LLFAT_KERNEL_EX(a, b)   LLFAT_CONCAT_EX(a,b)
 
#define PRECISION   1
 
#define Float   float
 
#define LOAD_FAT_MATRIX(gauge, dir, idx)   LOAD_MATRIX_18_SINGLE_DECLARE(gauge, dir, idx, FAT, fl.fat_ga_stride)
 
#define LOAD_EVEN_MULINK_MATRIX(dir, idx, var)   LOAD_MATRIX_18_SINGLE_TEX_DECLARE((odd_bit?muLink1TexSingle:muLink0TexSingle), dir, idx, var, fl.staple_stride)
 
#define LOAD_ODD_MULINK_MATRIX(dir, idx, var)   LOAD_MATRIX_18_SINGLE_TEX_DECLARE((odd_bit?muLink0TexSingle:muLink1TexSingle), dir, idx, var, fl.staple_stride)
 
#define LOAD_EVEN_FAT_MATRIX(dir, idx)   LOAD_MATRIX_18_SINGLE_TEX_DECLARE((odd_bit?fatGauge1TexSingle:fatGauge0TexSingle), dir, idx, FAT, fl.fat_ga_stride);
 
#define LOAD_ODD_FAT_MATRIX(dir, idx)   LOAD_MATRIX_18_SINGLE_TEX_DECLARE((odd_bit?fatGauge0TexSingle:fatGauge1TexSingle), dir, idx, FAT, fl.fat_ga_stride);
 
#define DECLARE_VAR_SIGN   short sign=1
 
#define SITELINK0TEX   siteLink0TexSingle_recon
 
#define SITELINK1TEX   siteLink1TexSingle_recon
 
#define LOAD_EVEN_SITE_MATRIX(dir, idx, var)   LOAD_MATRIX_12_SINGLE_DECLARE(sitelink_even, dir, idx, var, fl.site_ga_stride)
 
#define LOAD_ODD_SITE_MATRIX(dir, idx, var)   LOAD_MATRIX_12_SINGLE_DECLARE(sitelink_odd, dir, idx, var, fl.site_ga_stride)
 
#define LOAD_SITE_MATRIX(sitelink, dir, idx, var)   LOAD_MATRIX_12_SINGLE_DECLARE(sitelink, dir, idx, var, fl.site_ga_stride)
 
#define RECONSTRUCT_SITE_LINK(sign, var)   RECONSTRUCT_LINK_12(sign, var);
 
#define FloatN   float4
 
#define FloatM   float2
 
#define RECONSTRUCT   12
 
#define sd_data   float_12_sd_data
 
#define SITELINK0TEX   siteLink0TexSingle_norecon
 
#define SITELINK1TEX   siteLink1TexSingle_norecon
 
#define LOAD_EVEN_SITE_MATRIX(dir, idx, var)   LOAD_MATRIX_18_SINGLE_DECLARE(sitelink_even, dir, idx, var, fl.site_ga_stride)
 
#define LOAD_ODD_SITE_MATRIX(dir, idx, var)   LOAD_MATRIX_18_SINGLE_DECLARE(sitelink_odd, dir, idx, var, fl.site_ga_stride)
 
#define LOAD_SITE_MATRIX(sitelink, dir, idx, var)   LOAD_MATRIX_18_SINGLE_DECLARE(sitelink, dir, idx, var, fl.site_ga_stride)
 
#define RECONSTRUCT_SITE_LINK(sign, var)
 
#define FloatN   float2
 
#define FloatM   float2
 
#define RECONSTRUCT   18
 
#define sd_data   float_18_sd_data
 
#define PRECISION   0
 
#define Float   double
 
#define LOAD_FAT_MATRIX(gauge, dir, idx)   LOAD_MATRIX_18_DOUBLE_DECLARE(gauge, dir, idx, FAT, fl.fat_ga_stride)
 
#define LOAD_EVEN_MULINK_MATRIX(dir, idx, var)   LOAD_MATRIX_18_DOUBLE_TEX_DECLARE(odd_bit?muLink1TexDouble:muLink0TexDouble), mulink_even, dir, idx, var, fl.staple_stride)
 
#define LOAD_ODD_MULINK_MATRIX(dir, idx, var)   LOAD_MATRIX_18_DOUBLE_TEX_DECLARE((odd_bit?muLink0TexDouble:muLink1TexDouble), mulink_odd, dir, idx, var, fl.staple_stride)
 
#define LOAD_EVEN_FAT_MATRIX(dir, idx)   LOAD_MATRIX_18_DOUBLE_TEX_DECLARE((odd_bit?fatGauge1TexDouble:fatGauge0TexDouble), fatlink_even, dir, idx, FAT, fl.fat_ga_stride)
 
#define LOAD_ODD_FAT_MATRIX(dir, idx)   LOAD_MATRIX_18_DOUBLE_TEX_DECLARE((odd_bit?fatGauge0TexDouble:fatGauge1TexDouble), fatlink_odd, dir, idx, FAT, fl.fat_ga_stride)
 
#define SITELINK0TEX   siteLink0TexDouble
 
#define SITELINK1TEX   siteLink1TexDouble
 
#define LOAD_EVEN_SITE_MATRIX(dir, idx, var)   LOAD_MATRIX_18_DOUBLE_DECLARE(sitelink_even, dir, idx, var, fl.site_ga_stride)
 
#define LOAD_ODD_SITE_MATRIX(dir, idx, var)   LOAD_MATRIX_18_DOUBLE_DECLARE(sitelink_odd, dir, idx, var, fl.site_ga_stride)
 
#define LOAD_SITE_MATRIX(sitelink, dir, idx, var)   LOAD_MATRIX_18_DOUBLE_DECLARE(sitelink, dir, idx, var, fl.site_ga_stride)
 
#define RECONSTRUCT_SITE_LINK(sign, var)
 
#define FloatN   double2
 
#define FloatM   double2
 
#define RECONSTRUCT   18
 
#define sd_data   double_18_sd_data
 
#define SITELINK0TEX   siteLink0TexDouble
 
#define SITELINK1TEX   siteLink1TexDouble
 
#define LOAD_EVEN_SITE_MATRIX(dir, idx, var)   LOAD_MATRIX_12_DOUBLE_DECLARE(sitelink_even, dir, idx, var, fl.site_ga_stride)
 
#define LOAD_ODD_SITE_MATRIX(dir, idx, var)   LOAD_MATRIX_12_DOUBLE_DECLARE(sitelink_odd, dir, idx, var, fl.site_ga_stride)
 
#define LOAD_SITE_MATRIX(sitelink, dir, idx, var)   LOAD_MATRIX_12_DOUBLE_DECLARE(sitelink, dir, idx, var, fl.site_ga_stride)
 
#define RECONSTRUCT_SITE_LINK(sign, var)   RECONSTRUCT_LINK_12(sign, var);
 
#define FloatN   double2
 
#define FloatM   double2
 
#define RECONSTRUCT   12
 
#define sd_data   double_12_sd_data
 
#define UNBIND_ALL_TEXTURE
 
#define UNBIND_SITE_AND_FAT_LINK
 
#define BIND_MU_LINK()
 
#define UNBIND_MU_LINK()
 
#define BIND_SITE_AND_FAT_LINK
 
#define BIND_MU_LINK()
 
#define UNBIND_MU_LINK()
 
#define BIND_SITE_AND_FAT_LINK_REVERSE
 
#define ENUMERATE_FUNCS(mu, nu)
 
#define ENUMERATE_FUNCS_SAVE(mu, nu, save_staple)
 
#define CALL_FUNCTION(mu, nu)
 
#define CALL_FUNCTION(mu, nu, save_staple)
 
#define CALL_FUNCTION(mu, nu)
 
#define CALL_FUNCTION(mu, nu, save_staple)
 

Functions

void quda::llfat_init_cuda (QudaGaugeParam *param)
 
void quda::llfat_init_cuda_ex (QudaGaugeParam *param_ex)
 
template<int mu, int nu, int odd_bit>
__global__ void quda::LLFAT_KERNEL (do_siteComputeGenStapleParity, RECONSTRUCT)(FloatM *staple_even
 
 quda::if (kparam.kernel_type==LLFAT_EXTERIOR_KERNEL_FWD_X &&x1!=X1m1) return
 
 quda::COMPUTE_RECONSTRUCT_SIGN (sign, nu, x1, x2, x3, x4)
 
 quda::RECONSTRUCT_SITE_LINK (sign, a)
 
 quda::LLFAT_COMPUTE_NEW_IDX_PLUS (nu, X)
 
 quda::LOAD_ODD_SITE_MATRIX (mu, new_mem_idx, B)
 
 quda::COMPUTE_RECONSTRUCT_SIGN (sign, mu, new_x1, new_x2, new_x3, new_x4)
 
 quda::RECONSTRUCT_SITE_LINK (sign, b)
 
 quda::MULT_SU3_NN (a, b, tempa)
 
 quda::LLFAT_COMPUTE_NEW_IDX_PLUS (mu, X)
 
 quda::LOAD_ODD_SITE_MATRIX (nu, new_mem_idx, C)
 
 quda::COMPUTE_RECONSTRUCT_SIGN (sign, nu, new_x1, new_x2, new_x3, new_x4)
 
 quda::RECONSTRUCT_SITE_LINK (sign, c)
 
 quda::MULT_SU3_NA (tempa, c, staple)
 
 quda::LOAD_ODD_SITE_MATRIX (nu,(new_mem_idx), A)
 
 quda::LOAD_ODD_SITE_MATRIX (mu,(new_mem_idx), B)
 
 quda::MULT_SU3_AN (a, b, tempa)
 
 quda::LLFAT_COMPUTE_NEW_IDX_LOWER_STAPLE (nu, mu)
 
 quda::LOAD_EVEN_SITE_MATRIX (nu, new_mem_idx, C)
 
 quda::MULT_SU3_NN (tempa, c, b)
 
 quda::LLFAT_ADD_SU3_MATRIX (b, staple, staple)
 
 quda::WRITE_STAPLE_MATRIX (staple_even, mem_idx)
 
template<int mu, int nu, int odd_bit, int save_staple>
__global__ void quda::LLFAT_KERNEL (do_computeGenStapleFieldParity, RECONSTRUCT)(FloatM *staple_even
 
 quda::LOAD_ODD_MULINK_MATRIX (0, new_mem_idx, BB)
 
 quda::MULT_SU3_NN (a, bb, tempa)
 
 quda::if (save_staple)
 
 quda::LOAD_ODD_SITE_MATRIX (nu, new_mem_idx, A)
 
 quda::LLFAT_COMPUTE_NEW_IDX_MINUS (nu, X)
 
 quda::MULT_SU3_AN (a, bb, tempa)
 
 quda::MULT_SU3_NN (tempa, c, a)
 
 quda::LOAD_EVEN_FAT_MATRIX (mu, mem_idx)
 
__global__ void quda::LLFAT_KERNEL (llfatOneLink, RECONSTRUCT)(const FloatN *sitelink_even
 
 quda::if (mem_idx >=Vh)
 
 quda::for (int dir=0;dir< 4;dir++)
 
template<int mu, int nu, int odd_bit>
__global__ void quda::LLFAT_KERNEL_EX (do_siteComputeGenStapleParity, RECONSTRUCT)(FloatM *staple_even
 
 quda::if (mem_idx >=kparam.threads) return
 
 quda::COMPUTE_RECONSTRUCT_SIGN (sign, nu,(x1-2),(x2-2),(x3-2),(x4-2))
 
 quda::LLFAT_COMPUTE_NEW_IDX_PLUS_EX (nu, X)
 
 quda::COMPUTE_RECONSTRUCT_SIGN (sign, mu,(new_x1-2),(new_x2-2),(new_x3-2),(new_x4-2))
 
 quda::LLFAT_COMPUTE_NEW_IDX_PLUS_EX (mu, X)
 
 quda::COMPUTE_RECONSTRUCT_SIGN (sign, nu,(new_x1-2),(new_x2-2),(new_x3-2),(new_x4-2))
 
 quda::LLFAT_COMPUTE_NEW_IDX_LOWER_STAPLE_EX (nu, mu)
 
 quda::if (!(x1==1||x1==X1+2||x2==1||x2==X2+2||x3==1||x3==X3+2||x4==1||x4==X4+2))
 
template<int mu, int nu, int odd_bit, int save_staple>
__global__ void quda::LLFAT_KERNEL_EX (do_computeGenStapleFieldParity, RECONSTRUCT)(FloatM *staple_even
 
 quda::LLFAT_COMPUTE_NEW_IDX_MINUS_EX (nu, X)
 
 quda::LLFAT_ADD_SU3_MATRIX (a, staple, staple)
 
__global__ void quda::LLFAT_KERNEL_EX (llfatOneLink, RECONSTRUCT)(const FloatN *sitelink_even
 
 quda::if (sid >=2 *kparam.threads) return
 
 quda::if (idx >=kparam.threads)
 
void quda::siteComputeGenStapleParityKernel (void *staple_even, void *staple_odd, const void *sitelink_even, const void *sitelink_odd, void *fatlink_even, void *fatlink_odd, int mu, int nu, double mycoeff, QudaReconstructType recon, QudaPrecision prec, dim3 halfGridDim, llfat_kernel_param_t kparam, cudaStream_t *stream)
 
void quda::computeGenStapleFieldParityKernel (void *staple_even, void *staple_odd, const void *sitelink_even, const void *sitelink_odd, void *fatlink_even, void *fatlink_odd, const void *mulink_even, const void *mulink_odd, int mu, int nu, int save_staple, double mycoeff, QudaReconstructType recon, QudaPrecision prec, dim3 halfGridDim, llfat_kernel_param_t kparam, cudaStream_t *stream)
 
void quda::siteComputeGenStapleParityKernel_ex (void *staple_even, void *staple_odd, const void *sitelink_even, const void *sitelink_odd, void *fatlink_even, void *fatlink_odd, int mu, int nu, double mycoeff, QudaReconstructType recon, QudaPrecision prec, llfat_kernel_param_t kparam)
 
void quda::computeGenStapleFieldParityKernel_ex (void *staple_even, void *staple_odd, const void *sitelink_even, const void *sitelink_odd, void *fatlink_even, void *fatlink_odd, const void *mulink_even, const void *mulink_odd, int mu, int nu, int save_staple, double mycoeff, QudaReconstructType recon, QudaPrecision prec, llfat_kernel_param_t kparam)
 
void quda::llfatOneLinkKernel (cudaGaugeField &cudaFatLink, cudaGaugeField &cudaSiteLink, cudaGaugeField &cudaStaple, cudaGaugeField &cudaStaple1, QudaGaugeParam *param, double *act_path_coeff)
 
void quda::llfatOneLinkKernel_ex (cudaGaugeField &cudaFatLink, cudaGaugeField &cudaSiteLink, cudaGaugeField &cudaStaple, cudaGaugeField &cudaStaple1, QudaGaugeParam *param, double *act_path_coeff, llfat_kernel_param_t kparam)
 

Variables

__constant__ int quda::dir1_array [16]
 
__constant__ int quda::dir2_array [16]
 
unsigned long quda::staple_bytes =0
 
__global__ void FloatM * quda::staple_odd
 
__global__ void FloatM const
FloatN * 		quda::sitelink_even
 
__global__ void FloatM const
FloatN const FloatN * 		quda::sitelink_odd
 
__global__ void FloatM const
FloatN const FloatN FloatM * 		quda::fatlink_even
 
__global__ void FloatM const
FloatN const FloatN FloatM
FloatM * 		quda::fatlink_odd
 
__global__ void FloatM const
FloatN const FloatN FloatM
FloatM Float 		quda::mycoeff
 
__global__ void FloatM const
FloatN const FloatN FloatM
FloatM Float
llfat_kernel_param_t 		quda::kparam
 
FloatM quda::TEMPA5
 
FloatM quda::TEMPA6
 
FloatM quda::TEMPA7
 
FloatM quda::TEMPA8
 
FloatM quda::STAPLE0
 
FloatM quda::STAPLE1
 
FloatM quda::STAPLE2
 
FloatM quda::STAPLE3
 
FloatM quda::STAPLE4
 
FloatM quda::STAPLE5
 
FloatM quda::STAPLE6
 
FloatM quda::STAPLE7
 
FloatM quda::STAPLE8
 
int quda::mem_idx = blockIdx.x*blockDim.x + threadIdx.x
 
int quda::z1 = mem_idx / X1h
 
short quda::x1h = mem_idx - z1*X1h
 
int quda::z2 = z1 / X2
 
short quda::x2 = z1 - z2*X2
 
short quda::x4 = z2 / X3
 
short quda::x3 = z2 - x4*X3
 
short quda::x1odd = (x2 + x3 + x4 + odd_bit) & 1
 
short quda::x1 = 2*x1h + x1odd
 
int quda::X = 2*mem_idx + x1odd
 
int quda::new_mem_idx
 
 quda::DECLARE_VAR_SIGN
 
 quda::DECLARE_NEW_X
 
 quda::DECLARE_X_ARRAY
 
 quda::return
 
__global__ void FloatM const
FloatN const FloatN FloatM
FloatM const FloatM * 		quda::mulink_even
 
__global__ void FloatM const
FloatN const FloatN FloatM
FloatM const FloatM const
FloatM * 		quda::mulink_odd
 
FloatM quda::TEMPB0
 
FloatM quda::TEMPB1
 
FloatM quda::TEMPB2
 
FloatM quda::TEMPB3
 
FloatM quda::TEMPB4
 
FloatM quda::TEMPB5
 
FloatM quda::TEMPB6
 
FloatM quda::TEMPB7
 
FloatM quda::TEMPB8
 
 quda::else
 
__global__ void const FloatN
FloatM FloatM Float 		quda::coeff0
 
__global__ void const FloatN
FloatM FloatM Float Float 		quda::coeff5
 
FloatM * quda::my_fatlink = fatlink_even
 
int quda::sid = blockIdx.x*blockDim.x + threadIdx.x
 
int quda::odd_bit = 0
 
 quda::my_sitelink = sitelink_even
 
int quda::idx = sid
 

Macro Definition Documentation

#define BIND_MU_LINK ( )
Value:
do{ \
if(prec == QUDA_DOUBLE_PRECISION){ \
cudaBindTexture(0, muLink0TexDouble, mulink_even, staple_bytes); \
cudaBindTexture(0, muLink1TexDouble, mulink_odd, staple_bytes); \
}else{ \
cudaBindTexture(0, muLink0TexSingle, mulink_even, staple_bytes); \
cudaBindTexture(0, muLink1TexSingle, mulink_odd, staple_bytes); \
} \
}while(0)

Definition at line 539 of file llfat_quda.cu.

#define BIND_MU_LINK ( )
Value:
do{ \
if(prec == QUDA_DOUBLE_PRECISION){ \
cudaBindTexture(0, muLink0TexDouble, mulink_even, staple_bytes); \
cudaBindTexture(0, muLink1TexDouble, mulink_odd, staple_bytes); \
}else{ \
cudaBindTexture(0, muLink0TexSingle, mulink_even, staple_bytes); \
cudaBindTexture(0, muLink1TexSingle, mulink_odd, staple_bytes); \
} \
}while(0)

Definition at line 539 of file llfat_quda.cu.

#define BIND_SITE_AND_FAT_LINK
Value:
do { \
if(prec == QUDA_DOUBLE_PRECISION){ \
cudaBindTexture(0, siteLink0TexDouble, cudaSiteLink.Even_p(), cudaSiteLink.Bytes()); \
cudaBindTexture(0, siteLink1TexDouble, cudaSiteLink.Odd_p(), cudaSiteLink.Bytes()); \
cudaBindTexture(0, fatGauge0TexDouble, cudaFatLink.Even_p(), cudaFatLink.Bytes()); \
cudaBindTexture(0, fatGauge1TexDouble, cudaFatLink.Odd_p(), cudaFatLink.Bytes()); \
}else{ \
if(cudaSiteLink.Reconstruct() == QUDA_RECONSTRUCT_NO){ \
cudaBindTexture(0, siteLink0TexSingle_norecon, cudaSiteLink.Even_p(), cudaSiteLink.Bytes()); \
cudaBindTexture(0, siteLink1TexSingle_norecon, cudaSiteLink.Odd_p(), cudaSiteLink.Bytes()); \
}else{ \
cudaBindTexture(0, siteLink0TexSingle_recon, cudaSiteLink.Even_p(), cudaSiteLink.Bytes()); \
cudaBindTexture(0, siteLink1TexSingle_recon, cudaSiteLink.Odd_p(), cudaSiteLink.Bytes()); \
} \
cudaBindTexture(0, fatGauge0TexSingle, cudaFatLink.Even_p(), cudaFatLink.Bytes()); \
cudaBindTexture(0, fatGauge1TexSingle, cudaFatLink.Odd_p(), cudaFatLink.Bytes()); \
} \
}while(0)

Definition at line 520 of file llfat_quda.cu.

#define BIND_SITE_AND_FAT_LINK_REVERSE
Value:
do { \
if(prec == QUDA_DOUBLE_PRECISION){ \
cudaBindTexture(0, siteLink1TexDouble, cudaSiteLink.even, cudaSiteLink.bytes); \
cudaBindTexture(0, siteLink0TexDouble, cudaSiteLink.odd, cudaSiteLink.bytes); \
cudaBindTexture(0, fatGauge1TexDouble, cudaFatLink.even, cudaFatLink.bytes); \
cudaBindTexture(0, fatGauge0TexDouble, cudaFatLink.odd, cudaFatLink.bytes); \
}else{ \
if(cudaSiteLink.reconstruct == QUDA_RECONSTRUCT_NO){ \
cudaBindTexture(0, siteLink1TexSingle_norecon, cudaSiteLink.even, cudaSiteLink.bytes); \
cudaBindTexture(0, siteLink0TexSingle_norecon, cudaSiteLink.odd, cudaSiteLink.bytes); \
}else{ \
cudaBindTexture(0, siteLink1TexSingle_recon, cudaSiteLink.even, cudaSiteLink.bytes); \
cudaBindTexture(0, siteLink0TexSingle_recon, cudaSiteLink.odd, cudaSiteLink.bytes); \
} \
cudaBindTexture(0, fatGauge1TexSingle, cudaFatLink.even, cudaFatLink.bytes); \
cudaBindTexture(0, fatGauge0TexSingle, cudaFatLink.odd, cudaFatLink.bytes); \
} \
}while(0)

Definition at line 559 of file llfat_quda.cu.

#define BLOCK_DIM   64

Definition at line 23 of file llfat_quda.cu.

#define CALL_FUNCTION (   mu,
  nu 
)
#define CALL_FUNCTION (   mu,
  nu,
  save_staple 
)
#define CALL_FUNCTION (   mu,
  nu 
)
#define CALL_FUNCTION (   mu,
  nu,
  save_staple 
)
#define DECLARE_VAR_SIGN   short sign=1

Definition at line 297 of file llfat_quda.cu.

#define ENUMERATE_FUNCS (   mu,
  nu 
)

Definition at line 580 of file llfat_quda.cu.

#define ENUMERATE_FUNCS_SAVE (   mu,
  nu,
  save_staple 
)

Definition at line 647 of file llfat_quda.cu.

#define FATLINK_LOAD_TEX   1

Definition at line 20 of file llfat_quda.cu.

#define Float   float

Definition at line 367 of file llfat_quda.cu.

#define Float   double

Definition at line 367 of file llfat_quda.cu.

#define FloatM   float2

Definition at line 429 of file llfat_quda.cu.

#define FloatM   float2

Definition at line 429 of file llfat_quda.cu.

#define FloatM   double2

Definition at line 429 of file llfat_quda.cu.

#define FloatM   double2

Definition at line 429 of file llfat_quda.cu.

#define FloatN   float4

Definition at line 428 of file llfat_quda.cu.

#define FloatN   float2

Definition at line 428 of file llfat_quda.cu.

#define FloatN   double2

Definition at line 428 of file llfat_quda.cu.

#define FloatN   double2

Definition at line 428 of file llfat_quda.cu.

#define LLFAT_ADD_SU3_MATRIX (   ma,
  mb,
  mc 
)
Value:
mc##00_re = ma##00_re + mb##00_re; \
mc##00_im = ma##00_im + mb##00_im; \
mc##01_re = ma##01_re + mb##01_re; \
mc##01_im = ma##01_im + mb##01_im; \
mc##02_re = ma##02_re + mb##02_re; \
mc##02_im = ma##02_im + mb##02_im; \
mc##10_re = ma##10_re + mb##10_re; \
mc##10_im = ma##10_im + mb##10_im; \
mc##11_re = ma##11_re + mb##11_re; \
mc##11_im = ma##11_im + mb##11_im; \
mc##12_re = ma##12_re + mb##12_re; \
mc##12_im = ma##12_im + mb##12_im; \
mc##20_re = ma##20_re + mb##20_re; \
mc##20_im = ma##20_im + mb##20_im; \
mc##21_re = ma##21_re + mb##21_re; \
mc##21_im = ma##21_im + mb##21_im; \
mc##22_re = ma##22_re + mb##22_re; \
mc##22_im = ma##22_im + mb##22_im;

Definition at line 168 of file llfat_quda.cu.

#define LLFAT_CONCAT (   a,
  b 
)    a##b##Kernel

Definition at line 268 of file llfat_quda.cu.

#define LLFAT_CONCAT_EX (   a,
  b 
)    a##b##Kernel_ex

Definition at line 269 of file llfat_quda.cu.

#define LLFAT_KERNEL (   a,
  b 
)    LLFAT_CONCAT(a,b)

Definition at line 270 of file llfat_quda.cu.

#define LLFAT_KERNEL_EX (   a,
  b 
)    LLFAT_CONCAT_EX(a,b)

Definition at line 271 of file llfat_quda.cu.

#define LOAD_EVEN_FAT_MATRIX (   dir,
  idx 
)    LOAD_MATRIX_18_SINGLE_TEX_DECLARE((odd_bit?fatGauge1TexSingle:fatGauge0TexSingle), dir, idx, FAT, fl.fat_ga_stride);

Definition at line 378 of file llfat_quda.cu.

#define LOAD_EVEN_FAT_MATRIX (   dir,
  idx 
)    LOAD_MATRIX_18_DOUBLE_TEX_DECLARE((odd_bit?fatGauge1TexDouble:fatGauge0TexDouble), fatlink_even, dir, idx, FAT, fl.fat_ga_stride)

Definition at line 378 of file llfat_quda.cu.

#define LOAD_EVEN_MULINK_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_18_SINGLE_TEX_DECLARE((odd_bit?muLink1TexSingle:muLink0TexSingle), dir, idx, var, fl.staple_stride)

Definition at line 370 of file llfat_quda.cu.

#define LOAD_EVEN_MULINK_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_18_DOUBLE_TEX_DECLARE(odd_bit?muLink1TexDouble:muLink0TexDouble), mulink_even, dir, idx, var, fl.staple_stride)

Definition at line 370 of file llfat_quda.cu.

#define LOAD_EVEN_SITE_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_12_SINGLE_DECLARE(sitelink_even, dir, idx, var, fl.site_ga_stride)

Definition at line 423 of file llfat_quda.cu.

#define LOAD_EVEN_SITE_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_18_SINGLE_DECLARE(sitelink_even, dir, idx, var, fl.site_ga_stride)

Definition at line 423 of file llfat_quda.cu.

#define LOAD_EVEN_SITE_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_18_DOUBLE_DECLARE(sitelink_even, dir, idx, var, fl.site_ga_stride)

Definition at line 423 of file llfat_quda.cu.

#define LOAD_EVEN_SITE_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_12_DOUBLE_DECLARE(sitelink_even, dir, idx, var, fl.site_ga_stride)

Definition at line 423 of file llfat_quda.cu.

#define LOAD_FAT_MATRIX (   gauge,
  dir,
  idx 
)    LOAD_MATRIX_18_SINGLE_DECLARE(gauge, dir, idx, FAT, fl.fat_ga_stride)

Definition at line 368 of file llfat_quda.cu.

#define LOAD_FAT_MATRIX (   gauge,
  dir,
  idx 
)    LOAD_MATRIX_18_DOUBLE_DECLARE(gauge, dir, idx, FAT, fl.fat_ga_stride)

Definition at line 368 of file llfat_quda.cu.

#define LOAD_MATRIX_12_DOUBLE_DECLARE (   gauge,
  dir,
  idx,
  var,
  stride 
)
Value:
double2 var##0 = gauge[idx + dir*6*stride]; \
double2 var##1 = gauge[idx + dir*6*stride + stride]; \
double2 var##2 = gauge[idx + dir*6*stride + 2*stride]; \
double2 var##3 = gauge[idx + dir*6*stride + 3*stride]; \
double2 var##4 = gauge[idx + dir*6*stride + 4*stride]; \
double2 var##5 = gauge[idx + dir*6*stride + 5*stride]; \
double2 var##6, var##7, var##8;

Definition at line 149 of file llfat_quda.cu.

#define LOAD_MATRIX_12_DOUBLE_TEX_DECLARE (   gauge_tex,
  gauge,
  dir,
  idx,
  var,
  stride 
)
Value:
double2 var##0 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*6*stride); \
double2 var##1 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*6*stride + stride); \
double2 var##2 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*6*stride + 2*stride); \
double2 var##3 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*6*stride + 3*stride); \
double2 var##4 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*6*stride + 4*stride); \
double2 var##5 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*6*stride + 5*stride); \
double2 var##6, var##7, var##8;

Definition at line 159 of file llfat_quda.cu.

#define LOAD_MATRIX_12_SINGLE_DECLARE (   gauge,
  dir,
  idx,
  var,
  stride 
)
Value:
float4 var##0 = gauge[idx + dir*3*stride]; \
float4 var##1 = gauge[idx + dir*3*stride + stride]; \
float4 var##2 = gauge[idx + dir*3*stride + 2*stride]; \
float4 var##3, var##4;

Definition at line 88 of file llfat_quda.cu.

#define LOAD_MATRIX_12_SINGLE_TEX_DECLARE (   gauge,
  dir,
  idx,
  var,
  stride 
)
Value:
float4 var##0 = tex1Dfetch(gauge, idx + dir*3*stride); \
float4 var##1 = tex1Dfetch(gauge, idx + dir*3*stride + stride); \
float4 var##2 = tex1Dfetch(gauge, idx + dir*3*stride + 2*stride); \
float4 var##3, var##4;

Definition at line 94 of file llfat_quda.cu.

#define LOAD_MATRIX_18_DOUBLE_DECLARE (   gauge,
  dir,
  idx,
  var,
  stride 
)
Value:
double2 var##0 = gauge[idx + dir*9*stride]; \
double2 var##1 = gauge[idx + dir*9*stride + stride]; \
double2 var##2 = gauge[idx + dir*9*stride + 2*stride]; \
double2 var##3 = gauge[idx + dir*9*stride + 3*stride]; \
double2 var##4 = gauge[idx + dir*9*stride + 4*stride]; \
double2 var##5 = gauge[idx + dir*9*stride + 5*stride]; \
double2 var##6 = gauge[idx + dir*9*stride + 6*stride]; \
double2 var##7 = gauge[idx + dir*9*stride + 7*stride]; \
double2 var##8 = gauge[idx + dir*9*stride + 8*stride];

Definition at line 125 of file llfat_quda.cu.

#define LOAD_MATRIX_18_DOUBLE_TEX_DECLARE (   gauge_tex,
  gauge,
  dir,
  idx,
  var,
  stride 
)
Value:
double2 var##0 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride); \
double2 var##1 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride + stride); \
double2 var##2 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride + 2*stride); \
double2 var##3 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride + 3*stride); \
double2 var##4 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride + 4*stride); \
double2 var##5 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride + 5*stride); \
double2 var##6 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride + 6*stride); \
double2 var##7 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride + 7*stride); \
double2 var##8 = READ_DOUBLE2_TEXTURE(gauge_tex, gauge, idx + dir*9*stride + 8*stride);

Definition at line 137 of file llfat_quda.cu.

#define LOAD_MATRIX_18_SINGLE_DECLARE (   gauge,
  dir,
  idx,
  var,
  stride 
)
Value:
float2 var##0 = gauge[idx + dir*9*stride]; \
float2 var##1 = gauge[idx + dir*9*stride + stride]; \
float2 var##2 = gauge[idx + dir*9*stride + 2*stride]; \
float2 var##3 = gauge[idx + dir*9*stride + 3*stride]; \
float2 var##4 = gauge[idx + dir*9*stride + 4*stride]; \
float2 var##5 = gauge[idx + dir*9*stride + 5*stride]; \
float2 var##6 = gauge[idx + dir*9*stride + 6*stride]; \
float2 var##7 = gauge[idx + dir*9*stride + 7*stride]; \
float2 var##8 = gauge[idx + dir*9*stride + 8*stride];

Definition at line 100 of file llfat_quda.cu.

#define LOAD_MATRIX_18_SINGLE_TEX_DECLARE (   gauge,
  dir,
  idx,
  var,
  stride 
)
Value:
float2 var##0 = tex1Dfetch(gauge, idx + dir*9*stride); \
float2 var##1 = tex1Dfetch(gauge, idx + dir*9*stride + stride); \
float2 var##2 = tex1Dfetch(gauge, idx + dir*9*stride + 2*stride); \
float2 var##3 = tex1Dfetch(gauge, idx + dir*9*stride + 3*stride); \
float2 var##4 = tex1Dfetch(gauge, idx + dir*9*stride + 4*stride); \
float2 var##5 = tex1Dfetch(gauge, idx + dir*9*stride + 5*stride); \
float2 var##6 = tex1Dfetch(gauge, idx + dir*9*stride + 6*stride); \
float2 var##7 = tex1Dfetch(gauge, idx + dir*9*stride + 7*stride); \
float2 var##8 = tex1Dfetch(gauge, idx + dir*9*stride + 8*stride);

Definition at line 112 of file llfat_quda.cu.

#define LOAD_ODD_FAT_MATRIX (   dir,
  idx 
)    LOAD_MATRIX_18_SINGLE_TEX_DECLARE((odd_bit?fatGauge0TexSingle:fatGauge1TexSingle), dir, idx, FAT, fl.fat_ga_stride);

Definition at line 379 of file llfat_quda.cu.

#define LOAD_ODD_FAT_MATRIX (   dir,
  idx 
)    LOAD_MATRIX_18_DOUBLE_TEX_DECLARE((odd_bit?fatGauge0TexDouble:fatGauge1TexDouble), fatlink_odd, dir, idx, FAT, fl.fat_ga_stride)

Definition at line 379 of file llfat_quda.cu.

#define LOAD_ODD_MULINK_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_18_SINGLE_TEX_DECLARE((odd_bit?muLink0TexSingle:muLink1TexSingle), dir, idx, var, fl.staple_stride)

Definition at line 371 of file llfat_quda.cu.

#define LOAD_ODD_MULINK_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_18_DOUBLE_TEX_DECLARE((odd_bit?muLink0TexDouble:muLink1TexDouble), mulink_odd, dir, idx, var, fl.staple_stride)

Definition at line 371 of file llfat_quda.cu.

#define LOAD_ODD_SITE_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_12_SINGLE_DECLARE(sitelink_odd, dir, idx, var, fl.site_ga_stride)

Definition at line 424 of file llfat_quda.cu.

#define LOAD_ODD_SITE_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_18_SINGLE_DECLARE(sitelink_odd, dir, idx, var, fl.site_ga_stride)

Definition at line 424 of file llfat_quda.cu.

#define LOAD_ODD_SITE_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_18_DOUBLE_DECLARE(sitelink_odd, dir, idx, var, fl.site_ga_stride)

Definition at line 424 of file llfat_quda.cu.

#define LOAD_ODD_SITE_MATRIX (   dir,
  idx,
  var 
)    LOAD_MATRIX_12_DOUBLE_DECLARE(sitelink_odd, dir, idx, var, fl.site_ga_stride)

Definition at line 424 of file llfat_quda.cu.

#define LOAD_SITE_MATRIX (   sitelink,
  dir,
  idx,
  var 
)    LOAD_MATRIX_12_SINGLE_DECLARE(sitelink, dir, idx, var, fl.site_ga_stride)

Definition at line 426 of file llfat_quda.cu.

#define LOAD_SITE_MATRIX (   sitelink,
  dir,
  idx,
  var 
)    LOAD_MATRIX_18_SINGLE_DECLARE(sitelink, dir, idx, var, fl.site_ga_stride)

Definition at line 426 of file llfat_quda.cu.

#define LOAD_SITE_MATRIX (   sitelink,
  dir,
  idx,
  var 
)    LOAD_MATRIX_18_DOUBLE_DECLARE(sitelink, dir, idx, var, fl.site_ga_stride)

Definition at line 426 of file llfat_quda.cu.

#define LOAD_SITE_MATRIX (   sitelink,
  dir,
  idx,
  var 
)    LOAD_MATRIX_12_DOUBLE_DECLARE(sitelink, dir, idx, var, fl.site_ga_stride)

Definition at line 426 of file llfat_quda.cu.

#define MULINK_LOAD_TEX   1

Definition at line 19 of file llfat_quda.cu.

#define PRECISION   1

Definition at line 366 of file llfat_quda.cu.

#define PRECISION   0

Definition at line 366 of file llfat_quda.cu.

#define RECONSTRUCT   12

Definition at line 430 of file llfat_quda.cu.

#define RECONSTRUCT   18

Definition at line 430 of file llfat_quda.cu.

#define RECONSTRUCT   18

Definition at line 430 of file llfat_quda.cu.

#define RECONSTRUCT   12

Definition at line 430 of file llfat_quda.cu.

#define RECONSTRUCT_SITE_LINK (   sign,
  var 
)    RECONSTRUCT_LINK_12(sign, var);

Definition at line 427 of file llfat_quda.cu.

#define RECONSTRUCT_SITE_LINK (   sign,
  var 
)

Definition at line 427 of file llfat_quda.cu.

#define RECONSTRUCT_SITE_LINK (   sign,
  var 
)

Definition at line 427 of file llfat_quda.cu.

#define RECONSTRUCT_SITE_LINK (   sign,
  var 
)    RECONSTRUCT_LINK_12(sign, var);

Definition at line 427 of file llfat_quda.cu.

#define SCALAR_MULT_SU3_MATRIX (   a,
  b,
  c 
)
Value:
c##00_re = a*b##00_re; \
c##00_im = a*b##00_im; \
c##01_re = a*b##01_re; \
c##01_im = a*b##01_im; \
c##02_re = a*b##02_re; \
c##02_im = a*b##02_im; \
c##10_re = a*b##10_re; \
c##10_im = a*b##10_im; \
c##11_re = a*b##11_re; \
c##11_im = a*b##11_im; \
c##12_re = a*b##12_re; \
c##12_im = a*b##12_im; \
c##20_re = a*b##20_re; \
c##20_im = a*b##20_im; \
c##21_re = a*b##21_re; \
c##21_im = a*b##21_im; \
c##22_re = a*b##22_re; \
c##22_im = a*b##22_im; \

Definition at line 49 of file llfat_quda.cu.

#define sd_data   float_12_sd_data

Definition at line 431 of file llfat_quda.cu.

#define sd_data   float_18_sd_data

Definition at line 431 of file llfat_quda.cu.

#define sd_data   double_18_sd_data

Definition at line 431 of file llfat_quda.cu.

#define sd_data   double_12_sd_data

Definition at line 431 of file llfat_quda.cu.

#define SITE_MATRIX_LOAD_TEX   0

Definition at line 18 of file llfat_quda.cu.

#define SITELINK0TEX   siteLink0TexSingle_recon

Definition at line 417 of file llfat_quda.cu.

#define SITELINK0TEX   siteLink0TexSingle_norecon

Definition at line 417 of file llfat_quda.cu.

#define SITELINK0TEX   siteLink0TexDouble

Definition at line 417 of file llfat_quda.cu.

#define SITELINK0TEX   siteLink0TexDouble

Definition at line 417 of file llfat_quda.cu.

#define SITELINK1TEX   siteLink1TexSingle_recon

Definition at line 418 of file llfat_quda.cu.

#define SITELINK1TEX   siteLink1TexSingle_norecon

Definition at line 418 of file llfat_quda.cu.

#define SITELINK1TEX   siteLink1TexDouble

Definition at line 418 of file llfat_quda.cu.

#define SITELINK1TEX   siteLink1TexDouble

Definition at line 418 of file llfat_quda.cu.

#define UNBIND_ALL_TEXTURE
Value:
do{ \
if(prec ==QUDA_DOUBLE_PRECISION){ \
cudaUnbindTexture(siteLink0TexDouble); \
cudaUnbindTexture(siteLink1TexDouble); \
cudaUnbindTexture(fatGauge0TexDouble); \
cudaUnbindTexture(fatGauge1TexDouble); \
cudaUnbindTexture(muLink0TexDouble); \
cudaUnbindTexture(muLink1TexDouble); \
}else{ \
if(cudaSiteLink.reconstruct == QUDA_RECONSTRUCT_NO){ \
cudaUnbindTexture(siteLink0TexSingle_norecon); \
cudaUnbindTexture(siteLink1TexSingle_norecon); \
}else{ \
cudaUnbindTexture(siteLink0TexSingle_recon); \
cudaUnbindTexture(siteLink1TexSingle_recon); \
} \
cudaUnbindTexture(fatGauge0TexSingle); \
cudaUnbindTexture(fatGauge1TexSingle); \
cudaUnbindTexture(muLink0TexSingle); \
cudaUnbindTexture(muLink1TexSingle); \
} \
}while(0)

Definition at line 456 of file llfat_quda.cu.

#define UNBIND_MU_LINK ( )
Value:
do{ \
if(prec == QUDA_DOUBLE_PRECISION){ \
cudaUnbindTexture(muLink0TexSingle); \
cudaUnbindTexture(muLink1TexSingle); \
}else{ \
cudaUnbindTexture(muLink0TexDouble); \
cudaUnbindTexture(muLink1TexDouble); \
} \
}while(0)

Definition at line 549 of file llfat_quda.cu.

#define UNBIND_MU_LINK ( )
Value:
do{ \
if(prec == QUDA_DOUBLE_PRECISION){ \
cudaUnbindTexture(muLink0TexSingle); \
cudaUnbindTexture(muLink1TexSingle); \
}else{ \
cudaUnbindTexture(muLink0TexDouble); \
cudaUnbindTexture(muLink1TexDouble); \
} \
}while(0)

Definition at line 549 of file llfat_quda.cu.

#define UNBIND_SITE_AND_FAT_LINK
Value:
do{ \
if(prec == QUDA_DOUBLE_PRECISION){ \
cudaUnbindTexture(siteLink0TexDouble); \
cudaUnbindTexture(siteLink1TexDouble); \
cudaUnbindTexture(fatGauge0TexDouble); \
cudaUnbindTexture(fatGauge1TexDouble); \
}else { \
if(cudaSiteLink.reconstruct == QUDA_RECONSTRUCT_NO){ \
cudaUnbindTexture(siteLink0TexSingle_norecon); \
cudaUnbindTexture(siteLink1TexSingle_norecon); \
}else{ \
cudaUnbindTexture(siteLink0TexSingle_recon); \
cudaUnbindTexture(siteLink1TexSingle_recon); \
} \
cudaUnbindTexture(fatGauge0TexSingle); \
cudaUnbindTexture(fatGauge1TexSingle); \
} \
}while(0)

Definition at line 479 of file llfat_quda.cu.

#define WRITE_FAT_MATRIX (   gauge,
  dir,
  idx 
)
Value:
do { \
gauge[idx + dir*9*fl.fat_ga_stride] = FAT0; \
gauge[idx + (dir*9+1) * fl.fat_ga_stride] = FAT1; \
gauge[idx + (dir*9+2) * fl.fat_ga_stride] = FAT2; \
gauge[idx + (dir*9+3) * fl.fat_ga_stride] = FAT3; \
gauge[idx + (dir*9+4) * fl.fat_ga_stride] = FAT4; \
gauge[idx + (dir*9+5) * fl.fat_ga_stride] = FAT5; \
gauge[idx + (dir*9+6) * fl.fat_ga_stride] = FAT6; \
gauge[idx + (dir*9+7) * fl.fat_ga_stride] = FAT7; \
gauge[idx + (dir*9+8) * fl.fat_ga_stride] = FAT8;} while(0)

Definition at line 25 of file llfat_quda.cu.

#define WRITE_STAPLE_MATRIX (   gauge,
  idx 
)
Value:
gauge[idx] = STAPLE0; \
gauge[idx + fl.staple_stride] = STAPLE1; \
gauge[idx + 2*fl.staple_stride] = STAPLE2; \
gauge[idx + 3*fl.staple_stride] = STAPLE3; \
gauge[idx + 4*fl.staple_stride] = STAPLE4; \
gauge[idx + 5*fl.staple_stride] = STAPLE5; \
gauge[idx + 6*fl.staple_stride] = STAPLE6; \
gauge[idx + 7*fl.staple_stride] = STAPLE7; \
gauge[idx + 8*fl.staple_stride] = STAPLE8;

Definition at line 37 of file llfat_quda.cu.

Generated on Wed Mar 20 2013 12:52:19 for QUDA by   doxygen 1.8.2