QUDA v0.4.0
A library for QCD on GPUs
Classes | Namespaces | Defines | Functions | Variables
quda/lib/hisq_paths_force_quda.cu File Reference
#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  hisq::fermion_force::PathCoefficients< Real >
struct  hisq::fermion_force::RealTypeId< float2 >
struct  hisq::fermion_force::RealTypeId< double2 >
struct  hisq::fermion_force::CoeffSign< pos_dir, odd_lattice >
struct  hisq::fermion_force::CoeffSign< 0, 1 >
struct  hisq::fermion_force::CoeffSign< 0, 0 >
struct  hisq::fermion_force::CoeffSign< 1, 1 >
struct  hisq::fermion_force::Sign< odd_lattice >
struct  hisq::fermion_force::Sign< 1 >
struct  hisq::fermion_force::ArrayLength< RealX >
struct  hisq::fermion_force::ArrayLength< float4 >

Namespaces

namespace  hisq
namespace  hisq::fermion_force

Defines

#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, Vh)
#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, Vh)
#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, Vh)
#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, Vh)
#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_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)
#define Pmu   tempmat[0]
#define P3   tempmat[1]
#define P5   tempmat[2]
#define Pnumu   tempmat[3]
#define Qmu   tempCmat[0]
#define Qnumu   tempCmat[1]

Functions

void hisq::fermion_force::hisqForceInitCuda (QudaGaugeParam *param)
__device__ float2 hisq::fermion_force::operator* (float a, const float2 &b)
__device__ double2 hisq::fermion_force::operator* (double a, const double2 &b)
__device__ const float2 & hisq::fermion_force::operator+= (float2 &a, const float2 &b)
__device__ const double2 & hisq::fermion_force::operator+= (double2 &a, const double2 &b)
__device__ const float4 & hisq::fermion_force::operator+= (float4 &a, const float4 &b)
template<class T >
__device__ void hisq::fermion_force::adjointMatrix (T *mat)
template<int N, class T >
__device__ void hisq::fermion_force::loadMatrixFromField (const T *const field_even, const T *const field_odd, int dir, int idx, T *const mat, int oddness)
template<class T >
__device__ void hisq::fermion_force::loadMatrixFromField (const T *const field_even, const T *const field_odd, int dir, int idx, T *const mat, int oddness)
__device__ void hisq::fermion_force::loadMatrixFromField (const float4 *const field_even, const float4 *const field_odd, int dir, int idx, float2 *const mat, int oddness)
template<class T >
__device__ void hisq::fermion_force::loadMatrixFromField (const T *const field_even, const T *const field_odd, int idx, T *const mat, int oddness)
template<class T , class U >
__device__ void hisq::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 hisq::fermion_force::addMatrixToField (const T *const mat, int idx, U coeff, T *const field_even, T *const field_odd, int oddness)
template<class T >
__device__ void hisq::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 hisq::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 hisq::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 hisq::fermion_force::reconstructSign (int *const sign, int dir, const T i[4])
template<class RealA , int oddBit>
__global__ void hisq::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)
template<class RealA , class RealB , int sig_positive, int mu_positive, int oddBit>
__global__ void hisq::fermion_force::HISQ_KERNEL_NAME (do_middle_link, EXT)(const RealA *const oprodEven
int sig_positive int
mu_positive int oddBit
__global__ void 		hisq::fermion_force::HISQ_KERNEL_NAME (do_side_link, EXT)(const RealA *const P3Even
short sig_positive short
mu_positive short oddBit
__global__ void 		hisq::fermion_force::HISQ_KERNEL_NAME (do_all_link, EXT)(const RealA *const oprodEven
int oddBit __global__ void hisq::fermion_force::HISQ_KERNEL_NAME (do_longlink, EXT)(const RealB *const linkEven
int oddBit __global__ void hisq::fermion_force::HISQ_KERNEL_NAME (do_complete_force, EXT)(const RealB *const linkEven
class RealB static void hisq::fermion_force::middle_link_kernel (const RealA *const oprodEven, const RealA *const oprodOdd, const RealA *const QprevEven, const RealA *const QprevOdd, const RealB *const linkEven, const RealB *const linkOdd, const cudaGaugeField &link, int sig, int mu, typename RealTypeId< RealA >::Type coeff, dim3 gridDim, dim3 BlockDim, RealA *const PmuEven, RealA *const PmuOdd, RealA *const P3Even, RealA *const P3Odd, RealA *const QmuEven, RealA *const QmuOdd, RealA *const newOprodEven, RealA *const newOprodOdd)
template<class RealA , class RealB >
void hisq::fermion_force::longlink_terms (const RealB *const linkEven, const RealB *const linkOdd, const RealA *const naikOprodEven, const RealA *const naikOprodOdd, int sig, typename RealTypeId< RealA >::Type naik_coeff, dim3 gridDim, dim3 blockDim, const cudaGaugeField &link, RealA *const outputEven, RealA *const outputOdd)
void hisq::fermion_force::hisqCompleteForceCuda (const QudaGaugeParam &param, const cudaGaugeField &oprod, const cudaGaugeField &link, cudaGaugeField *force)
void hisq::fermion_force::hisqLongLinkForceCuda (double coeff, const QudaGaugeParam &param, const cudaGaugeField &oprod, const cudaGaugeField &link, cudaGaugeField *newOprod)
void hisq::fermion_force::hisqStaplesForceCuda (const double path_coeff[6], const QudaGaugeParam &param, const cudaGaugeField &oprod, const cudaGaugeField &link, cudaGaugeField *newOprod)

Variables

texture< int4, 1 > hisq::fermion_force::newOprod0TexDouble
texture< int4, 1 > hisq::fermion_force::newOprod1TexDouble
texture< float2,
1, cudaReadModeElementType > 		hisq::fermion_force::newOprod0TexSingle
texture< float2,
1, cudaReadModeElementType > 		hisq::fermion_force::newOprod1TexSingle
__global__ void const RealA *const hisq::fermion_force::oprodOdd
__global__ void const RealA
*const const RealA *const 		hisq::fermion_force::QprevEven
__global__ void const RealA
*const const RealA *const
const RealA *const 		hisq::fermion_force::QprevOdd
__global__ void const RealA
*const const RealA *const
const RealA *const const RealB
*const 		hisq::fermion_force::linkEven
__global__ void const RealA
*const const RealA *const
const RealA *const const RealB
*const const RealB *const 		hisq::fermion_force::linkOdd
__global__ void const RealA
*const const RealA *const
const RealA *const const RealB
*const const RealB *const int 		hisq::fermion_force::sig
__global__ void const RealA
*const const RealA *const
const RealA *const const RealB
*const const RealB *const int
int 		hisq::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 		hisq::fermion_force::coeff
__global__ void const RealA
*const const RealA *const
const RealA *const const RealB
*const const RealB *const int
int RealTypeId< RealA >::Type
RealA *const 		hisq::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 		hisq::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 		hisq::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 		hisq::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 		hisq::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 		hisq::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 		hisq::fermion_force::newOprodEven
int sig_positive int
mu_positive int oddBit
__global__ void const RealA
*const const RealA *const 		hisq::fermion_force::QprodEven
int sig_positive int
mu_positive int oddBit
__global__ void const RealA
*const const RealA *const
const RealA *const 		hisq::fermion_force::QprodOdd
int sig_positive int
mu_positive int oddBit
__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 		hisq::fermion_force::accumu_coeff
int sig_positive int
mu_positive int oddBit
__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 		hisq::fermion_force::shortPEven
int sig_positive int
mu_positive int oddBit
__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 		hisq::fermion_force::shortPOdd
int oddBit __global__ void
const RealB *const const RealA
*const 		hisq::fermion_force::naikOprodEven
int oddBit __global__ void
const RealB *const const RealA
*const const RealA *const 		hisq::fermion_force::naikOprodOdd
int oddBit __global__ void
const RealB *const const RealA
*const const RealA *const int
RealTypeId< RealA >::Type
RealA *const 		hisq::fermion_force::outputEven
int oddBit __global__ void
const RealB *const const RealA
*const 		hisq::fermion_force::oprodEven
int oddBit __global__ void
const RealB *const const RealA
*const const RealA *const int
RealA *const 		hisq::fermion_force::forceEven

Define 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*Vhx9); \
      value[1] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*Vhx9 + Vh);     \
      value[2] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*Vhx9 + 2*Vh); \
      value[3] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*Vhx9 + 3*Vh); \
      value[4] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*Vhx9 + 4*Vh); \
      value[5] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*Vhx9 + 5*Vh); \
      value[6] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*Vhx9 + 6*Vh); \
      value[7] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*Vhx9 + 7*Vh); \
      value[8] = LOAD_TEX_ENTRY( ((oddness)?NEWOPROD_ODD_TEX:NEWOPROD_EVEN_TEX), field, idx+dir*Vhx9 + 8*Vh); \
      field[idx + dir*Vhx9]          = value[0] + coeff*mat[0];         \
      field[idx + dir*Vhx9 + Vh]     = value[1] + coeff*mat[1];         \
      field[idx + dir*Vhx9 + Vhx2]   = value[2] + coeff*mat[2];         \
      field[idx + dir*Vhx9 + Vhx3]   = value[3] + coeff*mat[3];         \
      field[idx + dir*Vhx9 + Vhx4]   = value[4] + coeff*mat[4];         \
      field[idx + dir*Vhx9 + Vhx5]   = value[5] + coeff*mat[5];         \
      field[idx + dir*Vhx9 + Vhx6]   = value[6] + coeff*mat[6];         \
      field[idx + dir*Vhx9 + Vhx7]   = value[7] + coeff*mat[7];         \
      field[idx + dir*Vhx9 + Vhx8]   = value[8] + coeff*mat[8];         \
  }while(0)

Definition at line 196 of file hisq_paths_force_quda.cu.

#define CALL_ALL_LINK_KERNEL (   sig_sign,
  mu_sign 
)
Value:
if(sizeof(RealA) == sizeof(float2)){                            \
        if(recon  == QUDA_RECONSTRUCT_NO){                              \
          do_all_link_sp_18_kernel<float2, float2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(float2, float2); \
          do_all_link_sp_18_kernel<float2, float2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(float2, float2); \
        }else{                                                          \
          do_all_link_sp_12_kernel<float2, float4, sig_sign, mu_sign, 0> CALL_ARGUMENTS(float2, float4); \
          do_all_link_sp_12_kernel<float2, float4, sig_sign, mu_sign, 1> CALL_ARGUMENTS(float2, float4); \
        }                                                               \
      }else{                                                            \
        if(recon  == QUDA_RECONSTRUCT_NO){                              \
          do_all_link_dp_18_kernel<double2, double2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(double2, double2); \
          do_all_link_dp_18_kernel<double2, double2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(double2, double2); \
        }else{                                                          \
          do_all_link_dp_12_kernel<double2, double2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(double2, double2); \
          do_all_link_dp_12_kernel<double2, double2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(double2, double2); \
        }                                                               \
      }
#define CALL_ARGUMENTS (   typeA,
  typeB 
)
Value:
<<<halfGridDim, blockDim>>>((typeA*)P3Even, (typeA*)P3Odd, \
                                                                    (typeA*)oprodEven,  (typeA*)oprodOdd, \
                                                                    (typeB*)linkEven, (typeB*)linkOdd, \
                                                                    sig, mu, \
                                                                    (typename RealTypeId<typeA>::Type) coeff, \
                                                                    (typename RealTypeId<typeA>::Type) accumu_coeff, \
                                                                    (typeA*)shortPEven, (typeA*)shortPOdd, \
                                                                    (typeA*)newOprodEven, (typeA*)newOprodOdd)
#define CALL_ARGUMENTS (   typeA,
  typeB 
)
Value:
<<<halfGridDim, BlockDim>>>((typeA*)oprodEven, (typeA*)oprodOdd, \
                                                                 (typeA*)QprevEven, (typeA*)QprevOdd, \
                                                                 (typeB*)linkEven, (typeB*)linkOdd, \
                                                                 sig, mu, (typename RealTypeId<typeA>::Type)coeff, \
                                                                 (typeA*)PmuEven, (typeA*)PmuOdd, \
                                                                 (typeA*)P3Even, (typeA*)P3Odd, \
                                                                 (typeA*)QmuEven, (typeA*)QmuOdd, \
                                                                 (typeA*)newOprodEven, (typeA*)newOprodOdd)
#define CALL_ARGUMENTS (   typeA,
  typeB 
)
Value:
<<<halfGridDim, blockDim>>>((typeA*)oprodEven, (typeA*)oprodOdd, \
                                                                 (typeA*)QprevEven, (typeA*)QprevOdd, \
                                                                 (typeB*)linkEven, (typeB*)linkOdd, \
                                                                 sig,  mu, \
                                                                 (typename RealTypeId<typeA>::Type)coeff, \
                                                                 (typename RealTypeId<typeA>::Type)accumu_coeff, \
                                                                 (typeA*)shortPEven,(typeA*)shortPOdd, \
                                                                 (typeA*)newOprodEven, (typeA*)newOprodOdd)
#define CALL_ARGUMENTS (   typeA,
  typeB 
)
Value:
<<<halfGridDim,blockDim>>>((typeB*)linkEven, (typeB*)linkOdd, \
                                                                   (typeA*)naikOprodEven,  (typeA*)naikOprodOdd, \
                                                                   sig, naik_coeff, \
                                                                   (typeA*)outputEven, (typeA*)outputOdd); \
#define CALL_ARGUMENTS (   typeA,
  typeB 
)
Value:
<<<halfGridDim, blockDim>>>((typeB*)linkEven, (typeB*)linkOdd, \
                                                                  (typeA*)oprodEven, (typeA*)oprodOdd, \
                                                                  sig,  \
                                                                  (typeA*)momEven, (typeA*)momOdd);
#define CALL_MIDDLE_LINK_KERNEL (   sig_sign,
  mu_sign 
)
Value:
if(sizeof(RealA) == sizeof(float2)){                            \
          if(recon  == QUDA_RECONSTRUCT_NO){                            \
            do_middle_link_sp_18_kernel<float2, float2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(float2, float2); \
            do_middle_link_sp_18_kernel<float2, float2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(float2, float2); \
          }else{                                                        \
            do_middle_link_sp_12_kernel<float2, float4, sig_sign, mu_sign, 0> CALL_ARGUMENTS(float2, float4); \
            do_middle_link_sp_12_kernel<float2, float4, sig_sign, mu_sign, 1> CALL_ARGUMENTS(float2, float4); \
          }                                                             \
        }else{                                                          \
          if(recon  == QUDA_RECONSTRUCT_NO){                            \
            do_middle_link_dp_18_kernel<double2, double2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(double2, double2); \
            do_middle_link_dp_18_kernel<double2, double2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(double2, double2); \
          }else{                                                        \
            do_middle_link_dp_12_kernel<double2, double2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(double2, double2); \
            do_middle_link_dp_12_kernel<double2, double2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(double2, double2); \
          }                                                             \
        }
#define CALL_SIDE_LINK_KERNEL (   sig_sign,
  mu_sign 
)
Value:
if(sizeof(RealA) == sizeof(float2)){                            \
        if(recon  == QUDA_RECONSTRUCT_NO){                              \
          do_side_link_sp_18_kernel<float2, float2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(float2, float2); \
          do_side_link_sp_18_kernel<float2, float2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(float2, float2); \
        }else{                                                          \
          do_side_link_sp_12_kernel<float2, float4, sig_sign, mu_sign, 0> CALL_ARGUMENTS(float2, float4); \
          do_side_link_sp_12_kernel<float2, float4, sig_sign, mu_sign, 1> CALL_ARGUMENTS(float2, float4); \
        }                                                               \
      }else{                                                            \
        if(recon  == QUDA_RECONSTRUCT_NO){                              \
          do_side_link_dp_18_kernel<double2, double2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(double2, double2); \
          do_side_link_dp_18_kernel<double2, double2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(double2, double2); \
        }else{                                                          \
          do_side_link_dp_12_kernel<double2, double2, sig_sign, mu_sign, 0> CALL_ARGUMENTS(double2, double2); \
          do_side_link_dp_12_kernel<double2, double2, sig_sign, mu_sign, 1> CALL_ARGUMENTS(double2, double2); \
        }                                                               \
      }
#define COMPILE_HISQ_DP_12

Definition at line 12 of file hisq_paths_force_quda.cu.

#define COMPILE_HISQ_DP_18

Definition at line 11 of file hisq_paths_force_quda.cu.

#define COMPILE_HISQ_SP_12

Definition at line 14 of file hisq_paths_force_quda.cu.

#define COMPILE_HISQ_SP_18

Definition at line 13 of file hisq_paths_force_quda.cu.

#define COMPUTE_LINK_SIGN (   sign,
  dir,
  x 
)

Definition at line 461 of file hisq_paths_force_quda.cu.

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

Definition at line 461 of file hisq_paths_force_quda.cu.

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

Definition at line 461 of file hisq_paths_force_quda.cu.

#define COMPUTE_LINK_SIGN (   sign,
  dir,
  x 
)

Definition at line 461 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 440 of file hisq_paths_force_quda.cu.

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

Definition at line 442 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, Vh)

Definition at line 457 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, Vh)

Definition at line 457 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, Vh)

Definition at line 457 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, Vh)

Definition at line 457 of file hisq_paths_force_quda.cu.

#define HISQ_NEW_OPROD_LOAD_TEX   1

Definition at line 18 of file hisq_paths_force_quda.cu.

#define HISQ_SITE_MATRIX_LOAD_TEX   1

Definition at line 17 of file hisq_paths_force_quda.cu.

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

Definition at line 448 of file hisq_paths_force_quda.cu.

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

Definition at line 448 of file hisq_paths_force_quda.cu.

#define NEWOPROD_EVEN_TEX   newOprod0TexDouble

Definition at line 445 of file hisq_paths_force_quda.cu.

#define NEWOPROD_EVEN_TEX   newOprod0TexSingle

Definition at line 445 of file hisq_paths_force_quda.cu.

#define NEWOPROD_ODD_TEX   newOprod1TexDouble

Definition at line 446 of file hisq_paths_force_quda.cu.

#define NEWOPROD_ODD_TEX   newOprod1TexSingle

Definition at line 446 of file hisq_paths_force_quda.cu.

#define P3   tempmat[1]
#define P5   tempmat[2]
#define Pmu   tempmat[0]
#define Pnumu   tempmat[3]
#define PRECISION   1

Definition at line 454 of file hisq_paths_force_quda.cu.

#define PRECISION   0

Definition at line 454 of file hisq_paths_force_quda.cu.

#define PRECISION   1

Definition at line 454 of file hisq_paths_force_quda.cu.

#define PRECISION   0

Definition at line 454 of file hisq_paths_force_quda.cu.

#define Qmu   tempCmat[0]
#define Qnumu   tempCmat[1]
#define RECON   18

Definition at line 455 of file hisq_paths_force_quda.cu.

#define RECON   12

Definition at line 455 of file hisq_paths_force_quda.cu.

#define RECON   18

Definition at line 455 of file hisq_paths_force_quda.cu.

#define RECON   12

Definition at line 455 of file hisq_paths_force_quda.cu.

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

Definition at line 462 of file hisq_paths_force_quda.cu.

#define RECONSTRUCT_SITE_LINK (   var,
  sign 
)

Definition at line 462 of file hisq_paths_force_quda.cu.

#define RECONSTRUCT_SITE_LINK (   var,
  sign 
)

Definition at line 462 of file hisq_paths_force_quda.cu.

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

Definition at line 462 of file hisq_paths_force_quda.cu.

 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Defines
Generated on Wed Apr 4 2012 21:08:07 for QUDA by  doxygen 1.7.4