QUDA  v0.7.0
A library for QCD on GPUs
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
Macros | Functions | Variables
hisq_paths_force_core.h File Reference

Go to the source code of this file.

Macros

#define EXT   _sp_12_
 
#define print_matrix(mul)
 

Functions

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
 
 getCoords (x, sid, kparam.D, oddBit)
 
 updateCoords (y, mymu,(mu_positive?-1:1), kparam.X, kparam.ghostDim[mymu])
 
 if (mu_positive)
 
 updateCoords (y, mysig,(sig_positive?1:-1), kparam.X, kparam.ghostDim[mysig])
 
 for (int dir=0;dir< 4;++dir) y[dir]
 
 if (!mu_positive)
 
 if (sig_positive)
 
 loadLink< 18 > (linkEven, linkOdd, mysig, ab_link_nbr_idx, Uab.data, sig_positive^(1-oddBit), kparam.thin_link_stride)
 
 loadLink< 18 > (linkEven, linkOdd, mymu, bc_link_nbr_idx, Ubc.data, mu_positive^(1-oddBit), kparam.thin_link_stride)
 
 if (QprevOdd==NULL)
 
 if (PmuOdd)
 
 storeMatrixToField (Oy.data, new_sid, P3Even, P3Odd, oddBit, kparam.color_matrix_stride)
 
 loadLink< 18 > (linkEven, linkOdd, mymu, ad_link_nbr_idx, Uad.data, mu_positive^oddBit, kparam.thin_link_stride)
 
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, Oy.data, oddBit, kparam.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, Oy.data, oddBit, kparam.color_matrix_stride)
 
 addMatrixToField (Ow.data, point_d, accumu_coeff, shortPEven, shortPOdd, 1-oddBit, kparam.color_matrix_stride)
 
 loadMatrixFromField (QprodEven, QprodOdd, point_d, Ox.data, 1-oddBit, kparam.color_matrix_stride)
 
 if (oddBit)
 
 addMatrixToNewOprod (Ow.data, OPP_DIR(mu), new_sid, mycoeff, newOprodEven, newOprodOdd, oddBit, kparam.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_short, EXT)(const RealA *const P3Even
 
template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>
__global__ void HISQ_KERNEL_NAME (do_all_link, EXT)(const RealA *const oprodEven
 
 updateCoords (y, mu, 1, kparam.X, kparam.ghostDim[mu])
 
 loadMatrixFromField (QprevEven, QprevOdd, point_d, Ox.data, 1-oddBit, kparam.color_matrix_stride)
 
 loadLink< 18 > (linkEven, linkOdd, mu, new_sid, Uad.data, oddBit, kparam.thin_link_stride)
 
 loadLink< 18 > (linkEven, linkOdd, mu, point_b, Ubc.data, 1-oddBit, kparam.thin_link_stride)
 
 loadLink< 18 > (linkEven, linkOdd, posDir(sig), ab_link_nbr_idx, Uab.data, sig_positive^(1-oddBit), kparam.thin_link_stride)
 
 addMatrixToNewOprod (Ow.data, mu, new_sid, Sign< _oddBit^oddness_change >::result *mycoeff, newOprodEven, newOprodOdd, oddBit, kparam.color_matrix_stride)
 
template<class RealA , class RealB , int oddBit>
__global__ void HISQ_KERNEL_NAME (do_longlink, EXT)(const RealB *const linkEven
 
template<class RealA , class RealB , int oddBit>
__global__ void HISQ_KERNEL_NAME (do_complete_force, EXT)(const RealB *const linkEven
 

Variables

__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
 
__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 dx [4] = {0,0,0,0}
 
int x [4]
 
Matrix< RealA, 3 > Uab
 
Matrix< RealA, 3 > Ubc
 
Matrix< RealA, 3 > Uad
 
Matrix< RealA, 3 > Ow = conj(Ox)*conj(Oy)
 
Matrix< RealA, 3 > Ox
 
Matrix< RealA, 3 > Oy
 
int point_b = linkIndex(y, dx, E)
 
int point_c = linkIndex(y, dx, E)
 
int point_d = linkIndex(y, dx, E)
 
int ad_link_nbr_idx
 
int ab_link_nbr_idx = (sig_positive) ? new_sid : point_b
 
int bc_link_nbr_idx
 
int mymu = posDir(mu)
 
int E [4] = {kparam.X[0], kparam.X[1], kparam.X[2], kparam.X[3]}
 
int new_sid = sid
 
int y [4] = {x[0], x[1], x[2], x[3]}
 
 else
 
int mysig = posDir(sig)
 
int oddBit = _oddBit
 
 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
 
Matrix< RealA, 3 > Oz = Ubc*Oy
 
__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 RealTypeId
< RealA >::Type RealA *const 
outputEven
 
__global__ void const RealB
*const const RealA *const
const RealA *const RealTypeId
< RealA >::Type RealA *const
RealA *const 
outputOdd
 
int point_a
 
int point_e
 
Matrix< RealA, 3 > Ude
 
Matrix< RealA, 3 > Uef
 
__global__ void const RealB
*const const RealA *const 
oprodEven
 
__global__ void const RealB
*const const RealA *const
const RealA *const RealA
*const 
forceEven
 
__global__ void const RealB
*const const RealA *const
const RealA *const RealA
*const RealA *const 
forceOdd
 

Macro Definition Documentation

#define EXT   _sp_12_

Definition at line 10 of file hisq_paths_force_core.h.

#define print_matrix (   mul)
Value:
printf(" (%f %f) (%f %f) (%f %f)\n", mul##00_re, mul##00_im, mul##01_re, mul##01_im, mul##02_re, mul##02_im); \
printf(" (%f %f) (%f %f) (%f %f)\n", mul##10_re, mul##10_im, mul##11_re, mul##11_im, mul##12_re, mul##12_im); \
printf(" (%f %f) (%f %f) (%f %f)\n", mul##20_re, mul##20_im, mul##21_re, mul##21_im, mul##22_re, mul##22_im);

Definition at line 14 of file hisq_paths_force_core.h.

Function Documentation

addMatrixToField ( Ow.  data,
point_d  ,
accumu_coeff  ,
shortPEven  ,
shortPOdd  ,
1-  oddBit,
kparam.  color_matrix_stride 
)
addMatrixToNewOprod ( Ow.  data,
OPP_DIR(mu ,
new_sid  ,
mycoeff  ,
newOprodEven  ,
newOprodOdd  ,
oddBit  ,
kparam.  color_matrix_stride 
)
addMatrixToNewOprod ( Ow.  data,
mu  ,
new_sid  ,
Sign< _oddBit^oddness_change >::result *  mycoeff,
newOprodEven  ,
newOprodOdd  ,
oddBit  ,
kparam.  color_matrix_stride 
)
for ( )

Definition at line 848 of file hisq_paths_force_core.h.

getCoords ( x  ,
sid  ,
kparam.  D,
oddBit   
)
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
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
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
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
template<class RealA , class RealB , int sig_positive, int mu_positive, int _oddBit, int oddness_change>
__global__ void HISQ_KERNEL_NAME ( do_all_link  ,
EXT   
) const
template<class RealA , class RealB , int oddBit>
__global__ void HISQ_KERNEL_NAME ( do_longlink  ,
EXT   
) const
template<class RealA , class RealB , int oddBit>
__global__ void HISQ_KERNEL_NAME ( do_complete_force  ,
EXT   
) const
if ( sid >=kparam.  threads)
if ( mu_positive  )

Definition at line 145 of file hisq_paths_force_core.h.

if ( mu_positive)

Definition at line 164 of file hisq_paths_force_core.h.

if ( sig_positive  )

Definition at line 168 of file hisq_paths_force_core.h.

if ( QprevOdd  = = NULL)

Definition at line 186 of file hisq_paths_force_core.h.

if ( PmuOdd  )

Definition at line 200 of file hisq_paths_force_core.h.

if ( oddBit  )

Definition at line 533 of file hisq_paths_force_core.h.

loadLink< 18 > ( linkEven  ,
linkOdd  ,
mysig  ,
ab_link_nbr_idx  ,
Uab.  data,
sig_positive^  1-oddBit,
kparam.  thin_link_stride 
)
loadLink< 18 > ( linkEven  ,
linkOdd  ,
mymu  ,
bc_link_nbr_idx  ,
Ubc.  data,
mu_positive^  1-oddBit,
kparam.  thin_link_stride 
)
loadLink< 18 > ( linkEven  ,
linkOdd  ,
mymu  ,
ad_link_nbr_idx  ,
Uad.  data,
mu_positive^  oddBit,
kparam.  thin_link_stride 
)
loadLink< 18 > ( linkEven  ,
linkOdd  ,
mu  ,
new_sid  ,
Uad.  data,
oddBit  ,
kparam.  thin_link_stride 
)
loadLink< 18 > ( linkEven  ,
linkOdd  ,
mu  ,
point_b  ,
Ubc.  data,
1-  oddBit,
kparam.  thin_link_stride 
)
loadLink< 18 > ( linkEven  ,
linkOdd  ,
posDir(sig ,
ab_link_nbr_idx  ,
Uab.  data,
sig_positive^  1-oddBit,
kparam.  thin_link_stride 
)
loadMatrixFromField ( oprodEven  ,
oprodOdd  ,
point_c  ,
Oy.  data,
oddBit  ,
kparam.  color_matrix_stride 
)
loadMatrixFromField ( P3Even  ,
P3Odd  ,
new_sid  ,
Oy.  data,
oddBit  ,
kparam.  color_matrix_stride 
)
loadMatrixFromField ( QprodEven  ,
QprodOdd  ,
point_d  ,
Ox.  data,
1-  oddBit,
kparam.  color_matrix_stride 
)
loadMatrixFromField ( QprevEven  ,
QprevOdd  ,
point_d  ,
Ox.  data,
1-  oddBit,
kparam.  color_matrix_stride 
)
storeMatrixToField ( Oy.  data,
new_sid  ,
P3Even  ,
P3Odd  ,
oddBit  ,
kparam.  color_matrix_stride 
)
updateCoords ( y  ,
mymu  ,
(mu_positive?-1:1)  ,
kparam.  X,
kparam.  ghostDim[mymu] 
)
updateCoords ( y  ,
mysig  ,
(sig_positive?1:-1)  ,
kparam.  X,
kparam.  ghostDim[mysig] 
)
updateCoords ( y  ,
mu  ,
,
kparam.  X,
kparam.  ghostDim[mu] 
)

Variable Documentation

int ab_link_nbr_idx = (sig_positive) ? new_sid : point_b

Definition at line 117 of file hisq_paths_force_core.h.

__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

Definition at line 435 of file hisq_paths_force_core.h.

int ad_link_nbr_idx

Definition at line 117 of file hisq_paths_force_core.h.

int bc_link_nbr_idx

Definition at line 117 of file hisq_paths_force_core.h.

__global__ void const RealB* const const RealA* const const RealA* const RealTypeId<RealA>::Type coeff

Definition at line 82 of file hisq_paths_force_core.h.

int dx = {0,0,0,0}

Definition at line 98 of file hisq_paths_force_core.h.

int E = {kparam.X[0], kparam.X[1], kparam.X[2], kparam.X[3]}

Definition at line 133 of file hisq_paths_force_core.h.

else
Initial value:

Definition at line 147 of file hisq_paths_force_core.h.

__global__ void const RealB* const const RealA* const const RealA* const RealA* const forceEven

Definition at line 885 of file hisq_paths_force_core.h.

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

Definition at line 885 of file hisq_paths_force_core.h.

__global__ void const FloatN *const const FloatN *const Float const llfat_kernel_param_t kparam
Initial value:
{
int oddBit = _oddBit
int oddBit

Definition at line 92 of file hisq_paths_force_core.h.

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

Definition at line 82 of file hisq_paths_force_core.h.

__global__ void const RealB *const linkOdd

Definition at line 82 of file hisq_paths_force_core.h.

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

Definition at line 82 of file hisq_paths_force_core.h.

__global__ void FloatM const FloatN const FloatN FloatM FloatM const FloatM const FloatM Float mycoeff = CoeffSign<sig_positive,_oddBit ^ oddness_change>::result*coeff

Definition at line 496 of file hisq_paths_force_core.h.

int mymu = posDir(mu)

Definition at line 118 of file hisq_paths_force_core.h.

int mysig = posDir(sig)

Definition at line 151 of file hisq_paths_force_core.h.

__global__ void const RealB* const const RealA* const naikOprodEven

Definition at line 803 of file hisq_paths_force_core.h.

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

Definition at line 803 of file hisq_paths_force_core.h.

int new_sid = sid

Definition at line 134 of file hisq_paths_force_core.h.

__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 RealA *const newOprodEven

Definition at line 82 of file hisq_paths_force_core.h.

__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 RealA *const RealA *const newOprodOdd

Definition at line 82 of file hisq_paths_force_core.h.

int oddBit = _oddBit

Definition at line 263 of file hisq_paths_force_core.h.

__global__ void const RealB* const const RealA* const oprodEven

Definition at line 885 of file hisq_paths_force_core.h.

__global__ void const RealB *const const RealA *const const RealA *const oprodOdd

Definition at line 82 of file hisq_paths_force_core.h.

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

Definition at line 803 of file hisq_paths_force_core.h.

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

Definition at line 803 of file hisq_paths_force_core.h.

Matrix< RealA, 3 > Ow = conj(Ox)*conj(Oy)

Definition at line 105 of file hisq_paths_force_core.h.

Matrix< RealA, 3 > Ox

Definition at line 105 of file hisq_paths_force_core.h.

Matrix< RealA, 3 > Oy

Definition at line 105 of file hisq_paths_force_core.h.

Matrix< RealA, 3 > Oz = Ubc*Oy

Definition at line 673 of file hisq_paths_force_core.h.

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

Definition at line 82 of file hisq_paths_force_core.h.

__global__ void const RealA *const P3Odd

Definition at line 82 of file hisq_paths_force_core.h.

__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

Definition at line 82 of file hisq_paths_force_core.h.

__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

Definition at line 82 of file hisq_paths_force_core.h.

int point_a

Definition at line 829 of file hisq_paths_force_core.h.

int point_b = linkIndex(y, dx, E)

Definition at line 116 of file hisq_paths_force_core.h.

const int & point_c = linkIndex(y, dx, E)

Definition at line 116 of file hisq_paths_force_core.h.

int point_d = linkIndex(y, dx, E)

Definition at line 116 of file hisq_paths_force_core.h.

int point_e

Definition at line 829 of file hisq_paths_force_core.h.

__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

Definition at line 82 of file hisq_paths_force_core.h.

__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

Definition at line 82 of file hisq_paths_force_core.h.

__global__ void const RealA *const const RealA *const QprevEven

Definition at line 82 of file hisq_paths_force_core.h.

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

Definition at line 82 of file hisq_paths_force_core.h.

__global__ void const RealA* const const RealA* const QprodEven

Definition at line 435 of file hisq_paths_force_core.h.

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

Definition at line 435 of file hisq_paths_force_core.h.

return

Definition at line 380 of file hisq_paths_force_core.h.

__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

Definition at line 435 of file hisq_paths_force_core.h.

__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

Definition at line 435 of file hisq_paths_force_core.h.

Definition at line 96 of file hisq_paths_force_core.h.

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

Definition at line 82 of file hisq_paths_force_core.h.

Matrix< RealA, 3 > Uab

Definition at line 104 of file hisq_paths_force_core.h.

Matrix< RealA, 3 > Uad

Definition at line 104 of file hisq_paths_force_core.h.

Matrix< RealA, 3 > Ubc

Definition at line 104 of file hisq_paths_force_core.h.

Matrix<RealA,3> Ude

Definition at line 844 of file hisq_paths_force_core.h.

Matrix<RealA,3> Uef

Definition at line 844 of file hisq_paths_force_core.h.

int x

Definition at line 99 of file hisq_paths_force_core.h.

int y[4] = {x[0], x[1], x[2], x[3]}

Definition at line 137 of file hisq_paths_force_core.h.