|
QUDA v0.4.0
A library for QCD on GPUs
|
|
QUDA v0.4.0
A library for QCD on GPUs
|
#include <read_gauge.h>#include <gauge_field.h>#include <fermion_force_quda.h>#include <force_common.h>#include <hw_quda.h>Go to the source code of this file.
Defines | |
| #define | LOAD_ANTI_HERMITIAN(src, dir, idx, var) LOAD_ANTI_HERMITIAN_DIRECT(src, dir, idx, var, Vh) |
| #define | LOAD_HW_SINGLE(hw_even, hw_odd, idx, var, oddness) |
| #define | WRITE_HW_SINGLE(hw_even, hw_odd, idx, var, oddness) |
| #define | LOAD_HW(hw_eve, hw_odd, idx, var, oddness) LOAD_HW_SINGLE(hw_eve, hw_odd, idx, var, oddness) |
| #define | WRITE_HW(hw_even, hw_odd, idx, var, oddness) WRITE_HW_SINGLE(hw_even, hw_odd, idx, var, oddness) |
| #define | LOAD_MATRIX(src, dir, idx, var) LOAD_MATRIX_12_SINGLE(src, dir, idx, var, Vh) |
| #define | FF_SITE_MATRIX_LOAD_TEX 1 |
| #define | linkEvenTex siteLink0TexSingle_recon |
| #define | linkOddTex siteLink1TexSingle_recon |
| #define | FF_LOAD_MATRIX(dir, idx, var, oddness) LOAD_MATRIX_12_SINGLE_TEX(((oddness)?linkOddTex:linkEvenTex), dir, idx, var, Vh) |
| #define | linka00_re LINKA0.x |
| #define | linka00_im LINKA0.y |
| #define | linka01_re LINKA0.z |
| #define | linka01_im LINKA0.w |
| #define | linka02_re LINKA1.x |
| #define | linka02_im LINKA1.y |
| #define | linka10_re LINKA1.z |
| #define | linka10_im LINKA1.w |
| #define | linka11_re LINKA2.x |
| #define | linka11_im LINKA2.y |
| #define | linka12_re LINKA2.z |
| #define | linka12_im LINKA2.w |
| #define | linka20_re LINKA3.x |
| #define | linka20_im LINKA3.y |
| #define | linka21_re LINKA3.z |
| #define | linka21_im LINKA3.w |
| #define | linka22_re LINKA4.x |
| #define | linka22_im LINKA4.y |
| #define | linkb00_re LINKB0.x |
| #define | linkb00_im LINKB0.y |
| #define | linkb01_re LINKB0.z |
| #define | linkb01_im LINKB0.w |
| #define | linkb02_re LINKB1.x |
| #define | linkb02_im LINKB1.y |
| #define | linkb10_re LINKB1.z |
| #define | linkb10_im LINKB1.w |
| #define | linkb11_re LINKB2.x |
| #define | linkb11_im LINKB2.y |
| #define | linkb12_re LINKB2.z |
| #define | linkb12_im LINKB2.w |
| #define | linkb20_re LINKB3.x |
| #define | linkb20_im LINKB3.y |
| #define | linkb21_re LINKB3.z |
| #define | linkb21_im LINKB3.w |
| #define | linkb22_re LINKB4.x |
| #define | linkb22_im LINKB4.y |
| #define | MAT_MUL_HW(M, HW, HWOUT) |
| #define | ADJ_MAT_MUL_HW(M, HW, HWOUT) |
| #define | SU3_PROJECTOR(va, vb, m) |
| #define | SCALAR_MULT_ADD_SU3_VECTOR(va, vb, s, vc) |
| #define | FF_COMPUTE_NEW_FULL_IDX_PLUS_UPDATE(mydir, idx, new_idx) |
| #define | FF_COMPUTE_NEW_FULL_IDX_MINUS_UPDATE(mydir, idx, new_idx) |
| #define | FF_COMPUTE_NEW_FULL_IDX_PLUS(old_x1, old_x2, old_x3, old_x4, idx, mydir, new_idx) |
| #define | FF_COMPUTE_NEW_FULL_IDX_MINUS(old_x1, old_x2, old_x3, old_x4, idx, mydir, new_idx) |
| #define | ADD_FORCE_TO_MOM(hw1, hw2, idx, dir, cf, oddness) |
| #define | FF_COMPUTE_RECONSTRUCT_SIGN(sign, dir, i1, i2, i3, i4) |
| #define | hwa00_re HWA0.x |
| #define | hwa00_im HWA0.y |
| #define | hwa01_re HWA1.x |
| #define | hwa01_im HWA1.y |
| #define | hwa02_re HWA2.x |
| #define | hwa02_im HWA2.y |
| #define | hwa10_re HWA3.x |
| #define | hwa10_im HWA3.y |
| #define | hwa11_re HWA4.x |
| #define | hwa11_im HWA4.y |
| #define | hwa12_re HWA5.x |
| #define | hwa12_im HWA5.y |
| #define | hwb00_re HWB0.x |
| #define | hwb00_im HWB0.y |
| #define | hwb01_re HWB1.x |
| #define | hwb01_im HWB1.y |
| #define | hwb02_re HWB2.x |
| #define | hwb02_im HWB2.y |
| #define | hwb10_re HWB3.x |
| #define | hwb10_im HWB3.y |
| #define | hwb11_re HWB4.x |
| #define | hwb11_im HWB4.y |
| #define | hwb12_re HWB5.x |
| #define | hwb12_im HWB5.y |
| #define | hwc00_re HWC0.x |
| #define | hwc00_im HWC0.y |
| #define | hwc01_re HWC1.x |
| #define | hwc01_im HWC1.y |
| #define | hwc02_re HWC2.x |
| #define | hwc02_im HWC2.y |
| #define | hwc10_re HWC3.x |
| #define | hwc10_im HWC3.y |
| #define | hwc11_re HWC4.x |
| #define | hwc11_im HWC4.y |
| #define | hwc12_re HWC5.x |
| #define | hwc12_im HWC5.y |
| #define | hwd00_re HWD0.x |
| #define | hwd00_im HWD0.y |
| #define | hwd01_re HWD1.x |
| #define | hwd01_im HWD1.y |
| #define | hwd02_re HWD2.x |
| #define | hwd02_im HWD2.y |
| #define | hwd10_re HWD3.x |
| #define | hwd10_im HWD3.y |
| #define | hwd11_re HWD4.x |
| #define | hwd11_im HWD4.y |
| #define | hwd12_re HWD5.x |
| #define | hwd12_im HWD5.y |
| #define | hwe00_re HWE0.x |
| #define | hwe00_im HWE0.y |
| #define | hwe01_re HWE1.x |
| #define | hwe01_im HWE1.y |
| #define | hwe02_re HWE2.x |
| #define | hwe02_im HWE2.y |
| #define | hwe10_re HWE3.x |
| #define | hwe10_im HWE3.y |
| #define | hwe11_re HWE4.x |
| #define | hwe11_im HWE4.y |
| #define | hwe12_re HWE5.x |
| #define | hwe12_im HWE5.y |
| #define | CALL_MIDDLE_LINK_KERNEL(sig_sign, mu_sign) |
| #define | CALL_SIDE_LINK_KERNEL(sig_sign, mu_sign) |
| #define | CALL_ALL_LINK_KERNEL(sig_sign, mu_sign) |
| #define | Pmu tempvec[0] |
| #define | Pnumu tempvec[1] |
| #define | Prhonumu tempvec[2] |
| #define | P7 tempvec[3] |
| #define | P7rho tempvec[4] |
| #define | P7rhonu tempvec[5] |
| #define | P5 tempvec[6] |
| #define | P3 tempvec[7] |
| #define | P5nu tempvec[3] |
| #define | P3mu tempvec[3] |
| #define | Popmu tempvec[4] |
| #define | Pmumumu tempvec[4] |
Functions | |
| void | fermion_force_init_cuda (QudaGaugeParam *param) |
| template<int sig_positive, int mu_positive, int oddBit, typename Float2 > | |
| __global__ void | do_middle_link_kernel (Float2 *tempxEven, Float2 *tempxOdd, Float2 *PmuEven, Float2 *PmuOdd, Float2 *P3Even, Float2 *P3Odd, int sig, int mu, Float2 coeff, float4 *linkEven, float4 *linkOdd, Float2 *momEven, Float2 *momOdd) |
| template<int sig_positive, int mu_positive, int oddBit, typename Float2 > | |
| __global__ void | do_side_link_kernel (Float2 *P3Even, Float2 *P3Odd, Float2 *P3muEven, Float2 *P3muOdd, Float2 *TempxEven, Float2 *TempxOdd, Float2 *PmuEven, Float2 *PmuOdd, Float2 *shortPEven, Float2 *shortPOdd, int sig, int mu, Float2 coeff, Float2 accumu_coeff, float4 *linkEven, float4 *linkOdd, Float2 *momEven, Float2 *momOdd) |
| template<int sig_positive, int mu_positive, int oddBit, typename Float2 > | |
| __global__ void | do_all_link_kernel (Float2 *tempxEven, Float2 *tempxOdd, Float2 *PmuEven, Float2 *PmuOdd, Float2 *P3Even, Float2 *P3Odd, Float2 *P3muEven, Float2 *P3muOdd, Float2 *shortPEven, Float2 *shortPOdd, int sig, int mu, Float2 coeff, Float2 mcoeff, Float2 accumu_coeff, float4 *linkEven, float4 *linkOdd, Float2 *momEven, Float2 *momOdd) |
| template<int oddBit, typename Float2 > | |
| __global__ void | do_one_and_naik_terms_kernel (Float2 *TempxEven, Float2 *TempxOdd, Float2 *PmuEven, Float2 *PmuOdd, Float2 *PnumuEven, Float2 *PnumuOdd, int mu, Float2 OneLink, Float2 Naik, Float2 mNaik, float4 *linkEven, float4 *linkOdd, Float2 *momEven, Float2 *momOdd) |
| void | fermion_force_cuda (double eps, double weight1, double weight2, void *act_path_coeff, FullHw cudaHw, cudaGaugeField &siteLink, cudaGaugeField &cudaMom, QudaGaugeParam *param) |
| #define ADD_FORCE_TO_MOM | ( | hw1, | |
| hw2, | |||
| idx, | |||
| dir, | |||
| cf, | |||
| oddness | |||
| ) |
do{ \ Float2 my_coeff; \ int mydir; \ if (GOES_BACKWARDS(dir)){ \ mydir=OPP_DIR(dir); \ my_coeff.x = -cf.x; \ my_coeff.y = -cf.y; \ }else{ \ mydir=dir; \ my_coeff.x = cf.x; \ my_coeff.y = cf.y; \ } \ Float2 tmp_coeff; \ tmp_coeff.x = my_coeff.x; \ tmp_coeff.y = my_coeff.y; \ if(oddness){ \ tmp_coeff.x = - my_coeff.x; \ tmp_coeff.y = - my_coeff.y; \ } \ Float2* mom = oddness?momOdd:momEven; \ LOAD_ANTI_HERMITIAN(mom, mydir, idx, AH); \ UNCOMPRESS_ANTI_HERMITIAN(ah, linka); \ SU3_PROJECTOR(hw1##0, hw2##0, linkb); \ SCALAR_MULT_ADD_SU3_MATRIX(linka, linkb, tmp_coeff.x, linka); \ SU3_PROJECTOR(hw1##1, hw2##1, linkb); \ SCALAR_MULT_ADD_SU3_MATRIX(linka, linkb, tmp_coeff.y, linka); \ MAKE_ANTI_HERMITIAN(linka, ah); \ WRITE_ANTI_HERMITIAN(mom, mydir, idx, AH, Vh); \ }while(0)
Definition at line 273 of file fermion_force_quda.cu.
| #define ADJ_MAT_MUL_HW | ( | M, | |
| HW, | |||
| HWOUT | |||
| ) |
Definition at line 125 of file fermion_force_quda.cu.
| #define CALL_ALL_LINK_KERNEL | ( | sig_sign, | |
| mu_sign | |||
| ) |
do_all_link_kernel<sig_sign,mu_sign,0><<<halfGridDim, blockDim>>>(tempxEven, tempxOdd, \
PmuEven, PmuOdd, \
P3Even, P3Odd, \
P3muEven, P3muOdd, \
shortPEven, shortPOdd, \
sig, mu, coeff, mcoeff, accumu_coeff, \
linkEven, linkOdd, \
momEven, momOdd); \
do_all_link_kernel<sig_sign,mu_sign,1><<<halfGridDim, blockDim>>>(tempxEven, tempxOdd, \
PmuEven, PmuOdd, \
P3Even, P3Odd, \
P3muEven, P3muOdd, \
shortPEven, shortPOdd, \
sig, mu, coeff, mcoeff, accumu_coeff, \
linkEven, linkOdd, \
momEven, momOdd);
| #define CALL_MIDDLE_LINK_KERNEL | ( | sig_sign, | |
| mu_sign | |||
| ) |
do_middle_link_kernel<sig_sign, mu_sign,0><<<halfGridDim, BlockDim>>>( tempxEven, tempxOdd, \
PmuEven, PmuOdd, \
P3Even, P3Odd, \
sig, mu, coeff, \
linkEven, linkOdd, \
momEven, momOdd); \
do_middle_link_kernel<sig_sign, mu_sign, 1><<<halfGridDim, BlockDim>>>(tempxEven, tempxOdd, \
PmuEven, PmuOdd, \
P3Even, P3Odd, \
sig, mu, coeff, \
linkEven, linkOdd, \
momEven, momOdd);
| #define CALL_SIDE_LINK_KERNEL | ( | sig_sign, | |
| mu_sign | |||
| ) |
do_side_link_kernel<sig_sign,mu_sign,0><<<halfGridDim, blockDim>>>( P3Even, P3Odd, \ P3muEven, P3muOdd, \ TempxEven, TempxOdd, \ PmuEven, PmuOdd, \ shortPEven, shortPOdd, \ sig, mu, coeff, accumu_coeff, \ linkEven, linkOdd, \ momEven, momOdd); \ do_side_link_kernel<sig_sign,mu_sign,1><<<halfGridDim, blockDim>>>( P3Even, P3Odd, \ P3muEven, P3muOdd, \ TempxEven, TempxOdd, \ PmuEven, PmuOdd, \ shortPEven, shortPOdd, \ sig, mu, coeff, accumu_coeff, \ linkEven, linkOdd, \ momEven, momOdd);
| #define FF_COMPUTE_NEW_FULL_IDX_MINUS | ( | old_x1, | |
| old_x2, | |||
| old_x3, | |||
| old_x4, | |||
| idx, | |||
| mydir, | |||
| new_idx | |||
| ) |
do { \ switch(mydir){ \ case 0: \ new_idx = ( (old_x1==0)?idx+X1m1:idx-1); \ break; \ case 1: \ new_idx = ( (old_x2==0)?idx+X2X1mX1:idx-X1); \ break; \ case 2: \ new_idx = ( (old_x3==0)?idx+X3X2X1mX2X1:idx-X2X1); \ break; \ case 3: \ new_idx = ( (old_x4==0)?idx+X4X3X2X1mX3X2X1:idx-X3X2X1); \ break; \ } \ }while(0)
Definition at line 255 of file fermion_force_quda.cu.
| #define FF_COMPUTE_NEW_FULL_IDX_MINUS_UPDATE | ( | mydir, | |
| idx, | |||
| new_idx | |||
| ) |
do { \ switch(mydir){ \ case 0: \ new_idx = ( (new_x1==0)?idx+X1m1:idx-1); \ new_x1 = (new_x1==0)?X1m1:new_x1 - 1; \ break; \ case 1: \ new_idx = ( (new_x2==0)?idx+X2X1mX1:idx-X1); \ new_x2 = (new_x2==0)?X2m1:new_x2 - 1; \ break; \ case 2: \ new_idx = ( (new_x3==0)?idx+X3X2X1mX2X1:idx-X2X1); \ new_x3 = (new_x3==0)?X3m1:new_x3 - 1; \ break; \ case 3: \ new_idx = ( (new_x4==0)?idx+X4X3X2X1mX3X2X1:idx-X3X2X1); \ new_x4 = (new_x4==0)?X4m1:new_x4 - 1; \ break; \ } \ }while(0)
Definition at line 216 of file fermion_force_quda.cu.
| #define FF_COMPUTE_NEW_FULL_IDX_PLUS | ( | old_x1, | |
| old_x2, | |||
| old_x3, | |||
| old_x4, | |||
| idx, | |||
| mydir, | |||
| new_idx | |||
| ) |
do { \ switch(mydir){ \ case 0: \ new_idx = ( (old_x1==X1m1)?idx-X1m1:idx+1); \ break; \ case 1: \ new_idx = ( (old_x2==X2m1)?idx-X2X1mX1:idx+X1); \ break; \ case 2: \ new_idx = ( (old_x3==X3m1)?idx-X3X2X1mX2X1:idx+X2X1); \ break; \ case 3: \ new_idx = ( (old_x4==X4m1)?idx-X4X3X2X1mX3X2X1:idx+X3X2X1); \ break; \ } \ }while(0)
Definition at line 238 of file fermion_force_quda.cu.
| #define FF_COMPUTE_NEW_FULL_IDX_PLUS_UPDATE | ( | mydir, | |
| idx, | |||
| new_idx | |||
| ) |
do { \ switch(mydir){ \ case 0: \ new_idx = ( (new_x1==X1m1)?idx-X1m1:idx+1); \ new_x1 = (new_x1==X1m1)?0:new_x1+1; \ break; \ case 1: \ new_idx = ( (new_x2==X2m1)?idx-X2X1mX1:idx+X1); \ new_x2 = (new_x2==X2m1)?0:new_x2+1; \ break; \ case 2: \ new_idx = ( (new_x3==X3m1)?idx-X3X2X1mX2X1:idx+X2X1); \ new_x3 = (new_x3==X3m1)?0:new_x3+1; \ break; \ case 3: \ new_idx = ( (new_x4==X4m1)?idx-X4X3X2X1mX3X2X1:idx+X3X2X1); \ new_x4 = (new_x4==X4m1)?0:new_x4+1; \ break; \ } \ }while(0)
Definition at line 195 of file fermion_force_quda.cu.
| #define FF_COMPUTE_RECONSTRUCT_SIGN | ( | sign, | |
| dir, | |||
| i1, | |||
| i2, | |||
| i3, | |||
| i4 | |||
| ) |
do { \ sign =1; \ switch(dir){ \ case XUP: \ if ( (i4 & 1) == 1){ \ sign = -1; \ } \ break; \ case YUP: \ if ( ((i4+i1) & 1) == 1){ \ sign = -1; \ } \ break; \ case ZUP: \ if ( ((i4+i1+i2) & 1) == 1){ \ sign = -1; \ } \ break; \ case TUP: \ if (i4 == X4m1 ){ \ sign = -1; \ } \ break; \ } \ }while (0)
Definition at line 304 of file fermion_force_quda.cu.
| #define FF_LOAD_MATRIX | ( | dir, | |
| idx, | |||
| var, | |||
| oddness | |||
| ) | LOAD_MATRIX_12_SINGLE_TEX(((oddness)?linkOddTex:linkEvenTex), dir, idx, var, Vh) |
Definition at line 41 of file fermion_force_quda.cu.
| #define FF_SITE_MATRIX_LOAD_TEX 1 |
Definition at line 35 of file fermion_force_quda.cu.
| #define hwa00_im HWA0.y |
Definition at line 332 of file fermion_force_quda.cu.
| #define hwa00_re HWA0.x |
Definition at line 331 of file fermion_force_quda.cu.
| #define hwa01_im HWA1.y |
Definition at line 334 of file fermion_force_quda.cu.
| #define hwa01_re HWA1.x |
Definition at line 333 of file fermion_force_quda.cu.
| #define hwa02_im HWA2.y |
Definition at line 336 of file fermion_force_quda.cu.
| #define hwa02_re HWA2.x |
Definition at line 335 of file fermion_force_quda.cu.
| #define hwa10_im HWA3.y |
Definition at line 338 of file fermion_force_quda.cu.
| #define hwa10_re HWA3.x |
Definition at line 337 of file fermion_force_quda.cu.
| #define hwa11_im HWA4.y |
Definition at line 340 of file fermion_force_quda.cu.
| #define hwa11_re HWA4.x |
Definition at line 339 of file fermion_force_quda.cu.
| #define hwa12_im HWA5.y |
Definition at line 342 of file fermion_force_quda.cu.
| #define hwa12_re HWA5.x |
Definition at line 341 of file fermion_force_quda.cu.
| #define hwb00_im HWB0.y |
Definition at line 345 of file fermion_force_quda.cu.
| #define hwb00_re HWB0.x |
Definition at line 344 of file fermion_force_quda.cu.
| #define hwb01_im HWB1.y |
Definition at line 347 of file fermion_force_quda.cu.
| #define hwb01_re HWB1.x |
Definition at line 346 of file fermion_force_quda.cu.
| #define hwb02_im HWB2.y |
Definition at line 349 of file fermion_force_quda.cu.
| #define hwb02_re HWB2.x |
Definition at line 348 of file fermion_force_quda.cu.
| #define hwb10_im HWB3.y |
Definition at line 351 of file fermion_force_quda.cu.
| #define hwb10_re HWB3.x |
Definition at line 350 of file fermion_force_quda.cu.
| #define hwb11_im HWB4.y |
Definition at line 353 of file fermion_force_quda.cu.
| #define hwb11_re HWB4.x |
Definition at line 352 of file fermion_force_quda.cu.
| #define hwb12_im HWB5.y |
Definition at line 355 of file fermion_force_quda.cu.
| #define hwb12_re HWB5.x |
Definition at line 354 of file fermion_force_quda.cu.
| #define hwc00_im HWC0.y |
Definition at line 358 of file fermion_force_quda.cu.
| #define hwc00_re HWC0.x |
Definition at line 357 of file fermion_force_quda.cu.
| #define hwc01_im HWC1.y |
Definition at line 360 of file fermion_force_quda.cu.
| #define hwc01_re HWC1.x |
Definition at line 359 of file fermion_force_quda.cu.
| #define hwc02_im HWC2.y |
Definition at line 362 of file fermion_force_quda.cu.
| #define hwc02_re HWC2.x |
Definition at line 361 of file fermion_force_quda.cu.
| #define hwc10_im HWC3.y |
Definition at line 364 of file fermion_force_quda.cu.
| #define hwc10_re HWC3.x |
Definition at line 363 of file fermion_force_quda.cu.
| #define hwc11_im HWC4.y |
Definition at line 366 of file fermion_force_quda.cu.
| #define hwc11_re HWC4.x |
Definition at line 365 of file fermion_force_quda.cu.
| #define hwc12_im HWC5.y |
Definition at line 368 of file fermion_force_quda.cu.
| #define hwc12_re HWC5.x |
Definition at line 367 of file fermion_force_quda.cu.
| #define hwd00_im HWD0.y |
Definition at line 371 of file fermion_force_quda.cu.
| #define hwd00_re HWD0.x |
Definition at line 370 of file fermion_force_quda.cu.
| #define hwd01_im HWD1.y |
Definition at line 373 of file fermion_force_quda.cu.
| #define hwd01_re HWD1.x |
Definition at line 372 of file fermion_force_quda.cu.
| #define hwd02_im HWD2.y |
Definition at line 375 of file fermion_force_quda.cu.
| #define hwd02_re HWD2.x |
Definition at line 374 of file fermion_force_quda.cu.
| #define hwd10_im HWD3.y |
Definition at line 377 of file fermion_force_quda.cu.
| #define hwd10_re HWD3.x |
Definition at line 376 of file fermion_force_quda.cu.
| #define hwd11_im HWD4.y |
Definition at line 379 of file fermion_force_quda.cu.
| #define hwd11_re HWD4.x |
Definition at line 378 of file fermion_force_quda.cu.
| #define hwd12_im HWD5.y |
Definition at line 381 of file fermion_force_quda.cu.
| #define hwd12_re HWD5.x |
Definition at line 380 of file fermion_force_quda.cu.
| #define hwe00_im HWE0.y |
Definition at line 384 of file fermion_force_quda.cu.
| #define hwe00_re HWE0.x |
Definition at line 383 of file fermion_force_quda.cu.
| #define hwe01_im HWE1.y |
Definition at line 386 of file fermion_force_quda.cu.
| #define hwe01_re HWE1.x |
Definition at line 385 of file fermion_force_quda.cu.
| #define hwe02_im HWE2.y |
Definition at line 388 of file fermion_force_quda.cu.
| #define hwe02_re HWE2.x |
Definition at line 387 of file fermion_force_quda.cu.
| #define hwe10_im HWE3.y |
Definition at line 390 of file fermion_force_quda.cu.
| #define hwe10_re HWE3.x |
Definition at line 389 of file fermion_force_quda.cu.
| #define hwe11_im HWE4.y |
Definition at line 392 of file fermion_force_quda.cu.
| #define hwe11_re HWE4.x |
Definition at line 391 of file fermion_force_quda.cu.
| #define hwe12_im HWE5.y |
Definition at line 394 of file fermion_force_quda.cu.
| #define hwe12_re HWE5.x |
Definition at line 393 of file fermion_force_quda.cu.
| #define linka00_im LINKA0.y |
Definition at line 48 of file fermion_force_quda.cu.
| #define linka00_re LINKA0.x |
Definition at line 47 of file fermion_force_quda.cu.
| #define linka01_im LINKA0.w |
Definition at line 50 of file fermion_force_quda.cu.
| #define linka01_re LINKA0.z |
Definition at line 49 of file fermion_force_quda.cu.
| #define linka02_im LINKA1.y |
Definition at line 52 of file fermion_force_quda.cu.
| #define linka02_re LINKA1.x |
Definition at line 51 of file fermion_force_quda.cu.
| #define linka10_im LINKA1.w |
Definition at line 54 of file fermion_force_quda.cu.
| #define linka10_re LINKA1.z |
Definition at line 53 of file fermion_force_quda.cu.
| #define linka11_im LINKA2.y |
Definition at line 56 of file fermion_force_quda.cu.
| #define linka11_re LINKA2.x |
Definition at line 55 of file fermion_force_quda.cu.
| #define linka12_im LINKA2.w |
Definition at line 58 of file fermion_force_quda.cu.
| #define linka12_re LINKA2.z |
Definition at line 57 of file fermion_force_quda.cu.
| #define linka20_im LINKA3.y |
Definition at line 60 of file fermion_force_quda.cu.
| #define linka20_re LINKA3.x |
Definition at line 59 of file fermion_force_quda.cu.
| #define linka21_im LINKA3.w |
Definition at line 62 of file fermion_force_quda.cu.
| #define linka21_re LINKA3.z |
Definition at line 61 of file fermion_force_quda.cu.
| #define linka22_im LINKA4.y |
Definition at line 64 of file fermion_force_quda.cu.
| #define linka22_re LINKA4.x |
Definition at line 63 of file fermion_force_quda.cu.
| #define linkb00_im LINKB0.y |
Definition at line 67 of file fermion_force_quda.cu.
| #define linkb00_re LINKB0.x |
Definition at line 66 of file fermion_force_quda.cu.
| #define linkb01_im LINKB0.w |
Definition at line 69 of file fermion_force_quda.cu.
| #define linkb01_re LINKB0.z |
Definition at line 68 of file fermion_force_quda.cu.
| #define linkb02_im LINKB1.y |
Definition at line 71 of file fermion_force_quda.cu.
| #define linkb02_re LINKB1.x |
Definition at line 70 of file fermion_force_quda.cu.
| #define linkb10_im LINKB1.w |
Definition at line 73 of file fermion_force_quda.cu.
| #define linkb10_re LINKB1.z |
Definition at line 72 of file fermion_force_quda.cu.
| #define linkb11_im LINKB2.y |
Definition at line 75 of file fermion_force_quda.cu.
| #define linkb11_re LINKB2.x |
Definition at line 74 of file fermion_force_quda.cu.
| #define linkb12_im LINKB2.w |
Definition at line 77 of file fermion_force_quda.cu.
| #define linkb12_re LINKB2.z |
Definition at line 76 of file fermion_force_quda.cu.
| #define linkb20_im LINKB3.y |
Definition at line 79 of file fermion_force_quda.cu.
| #define linkb20_re LINKB3.x |
Definition at line 78 of file fermion_force_quda.cu.
| #define linkb21_im LINKB3.w |
Definition at line 81 of file fermion_force_quda.cu.
| #define linkb21_re LINKB3.z |
Definition at line 80 of file fermion_force_quda.cu.
| #define linkb22_im LINKB4.y |
Definition at line 83 of file fermion_force_quda.cu.
| #define linkb22_re LINKB4.x |
Definition at line 82 of file fermion_force_quda.cu.
| #define linkEvenTex siteLink0TexSingle_recon |
Definition at line 37 of file fermion_force_quda.cu.
| #define linkOddTex siteLink1TexSingle_recon |
Definition at line 38 of file fermion_force_quda.cu.
| #define LOAD_ANTI_HERMITIAN | ( | src, | |
| dir, | |||
| idx, | |||
| var | |||
| ) | LOAD_ANTI_HERMITIAN_DIRECT(src, dir, idx, var, Vh) |
Definition at line 9 of file fermion_force_quda.cu.
| #define LOAD_HW | ( | hw_eve, | |
| hw_odd, | |||
| idx, | |||
| var, | |||
| oddness | |||
| ) | LOAD_HW_SINGLE(hw_eve, hw_odd, idx, var, oddness) |
Definition at line 31 of file fermion_force_quda.cu.
| #define LOAD_HW_SINGLE | ( | hw_even, | |
| hw_odd, | |||
| idx, | |||
| var, | |||
| oddness | |||
| ) |
| #define LOAD_MATRIX | ( | src, | |
| dir, | |||
| idx, | |||
| var | |||
| ) | LOAD_MATRIX_12_SINGLE(src, dir, idx, var, Vh) |
Definition at line 33 of file fermion_force_quda.cu.
| #define MAT_MUL_HW | ( | M, | |
| HW, | |||
| HWOUT | |||
| ) |
Definition at line 86 of file fermion_force_quda.cu.
| #define P3 tempvec[7] |
Definition at line 1184 of file fermion_force_quda.cu.
| #define P3mu tempvec[3] |
Definition at line 1186 of file fermion_force_quda.cu.
| #define P5 tempvec[6] |
Definition at line 1183 of file fermion_force_quda.cu.
| #define P5nu tempvec[3] |
Definition at line 1185 of file fermion_force_quda.cu.
| #define P7 tempvec[3] |
Definition at line 1180 of file fermion_force_quda.cu.
| #define P7rho tempvec[4] |
Definition at line 1181 of file fermion_force_quda.cu.
| #define P7rhonu tempvec[5] |
Definition at line 1182 of file fermion_force_quda.cu.
| #define Pmu tempvec[0] |
Definition at line 1177 of file fermion_force_quda.cu.
| #define Pmumumu tempvec[4] |
Definition at line 1188 of file fermion_force_quda.cu.
| #define Pnumu tempvec[1] |
Definition at line 1178 of file fermion_force_quda.cu.
| #define Popmu tempvec[4] |
Definition at line 1187 of file fermion_force_quda.cu.
| #define Prhonumu tempvec[2] |
Definition at line 1179 of file fermion_force_quda.cu.
| #define SCALAR_MULT_ADD_SU3_VECTOR | ( | va, | |
| vb, | |||
| s, | |||
| vc | |||
| ) |
do { \ vc##0_re = va##0_re + vb##0_re * s; \ vc##0_im = va##0_im + vb##0_im * s; \ vc##1_re = va##1_re + vb##1_re * s; \ vc##1_im = va##1_im + vb##1_im * s; \ vc##2_re = va##2_re + vb##2_re * s; \ vc##2_im = va##2_im + vb##2_im * s; \ }while (0)
Definition at line 185 of file fermion_force_quda.cu.
| #define SU3_PROJECTOR | ( | va, | |
| vb, | |||
| m | |||
| ) |
m##00_re = va##0_re * vb##0_re + va##0_im * vb##0_im; \ m##00_im = va##0_im * vb##0_re - va##0_re * vb##0_im; \ m##01_re = va##0_re * vb##1_re + va##0_im * vb##1_im; \ m##01_im = va##0_im * vb##1_re - va##0_re * vb##1_im; \ m##02_re = va##0_re * vb##2_re + va##0_im * vb##2_im; \ m##02_im = va##0_im * vb##2_re - va##0_re * vb##2_im; \ m##10_re = va##1_re * vb##0_re + va##1_im * vb##0_im; \ m##10_im = va##1_im * vb##0_re - va##1_re * vb##0_im; \ m##11_re = va##1_re * vb##1_re + va##1_im * vb##1_im; \ m##11_im = va##1_im * vb##1_re - va##1_re * vb##1_im; \ m##12_re = va##1_re * vb##2_re + va##1_im * vb##2_im; \ m##12_im = va##1_im * vb##2_re - va##1_re * vb##2_im; \ m##20_re = va##2_re * vb##0_re + va##2_im * vb##0_im; \ m##20_im = va##2_im * vb##0_re - va##2_re * vb##0_im; \ m##21_re = va##2_re * vb##1_re + va##2_im * vb##1_im; \ m##21_im = va##2_im * vb##1_re - va##2_re * vb##1_im; \ m##22_re = va##2_re * vb##2_re + va##2_im * vb##2_im; \ m##22_im = va##2_im * vb##2_re - va##2_re * vb##2_im;
Definition at line 164 of file fermion_force_quda.cu.
| #define WRITE_HW | ( | hw_even, | |
| hw_odd, | |||
| idx, | |||
| var, | |||
| oddness | |||
| ) | WRITE_HW_SINGLE(hw_even, hw_odd, idx, var, oddness) |
Definition at line 32 of file fermion_force_quda.cu.
| #define WRITE_HW_SINGLE | ( | hw_even, | |
| hw_odd, | |||
| idx, | |||
| var, | |||
| oddness | |||
| ) |
| __global__ void do_all_link_kernel | ( | Float2 * | tempxEven, |
| Float2 * | tempxOdd, | ||
| Float2 * | PmuEven, | ||
| Float2 * | PmuOdd, | ||
| Float2 * | P3Even, | ||
| Float2 * | P3Odd, | ||
| Float2 * | P3muEven, | ||
| Float2 * | P3muOdd, | ||
| Float2 * | shortPEven, | ||
| Float2 * | shortPOdd, | ||
| int | sig, | ||
| int | mu, | ||
| Float2 | coeff, | ||
| Float2 | mcoeff, | ||
| Float2 | accumu_coeff, | ||
| float4 * | linkEven, | ||
| float4 * | linkOdd, | ||
| Float2 * | momEven, | ||
| Float2 * | momOdd | ||
| ) |
Definition at line 807 of file fermion_force_quda.cu.
| __global__ void do_middle_link_kernel | ( | Float2 * | tempxEven, |
| Float2 * | tempxOdd, | ||
| Float2 * | PmuEven, | ||
| Float2 * | PmuOdd, | ||
| Float2 * | P3Even, | ||
| Float2 * | P3Odd, | ||
| int | sig, | ||
| int | mu, | ||
| Float2 | coeff, | ||
| float4 * | linkEven, | ||
| float4 * | linkOdd, | ||
| Float2 * | momEven, | ||
| Float2 * | momOdd | ||
| ) |
Definition at line 423 of file fermion_force_quda.cu.
| __global__ void do_one_and_naik_terms_kernel | ( | Float2 * | TempxEven, |
| Float2 * | TempxOdd, | ||
| Float2 * | PmuEven, | ||
| Float2 * | PmuOdd, | ||
| Float2 * | PnumuEven, | ||
| Float2 * | PnumuOdd, | ||
| int | mu, | ||
| Float2 | OneLink, | ||
| Float2 | Naik, | ||
| Float2 | mNaik, | ||
| float4 * | linkEven, | ||
| float4 * | linkOdd, | ||
| Float2 * | momEven, | ||
| Float2 * | momOdd | ||
| ) |
Definition at line 1061 of file fermion_force_quda.cu.
| __global__ void do_side_link_kernel | ( | Float2 * | P3Even, |
| Float2 * | P3Odd, | ||
| Float2 * | P3muEven, | ||
| Float2 * | P3muOdd, | ||
| Float2 * | TempxEven, | ||
| Float2 * | TempxOdd, | ||
| Float2 * | PmuEven, | ||
| Float2 * | PmuOdd, | ||
| Float2 * | shortPEven, | ||
| Float2 * | shortPOdd, | ||
| int | sig, | ||
| int | mu, | ||
| Float2 | coeff, | ||
| Float2 | accumu_coeff, | ||
| float4 * | linkEven, | ||
| float4 * | linkOdd, | ||
| Float2 * | momEven, | ||
| Float2 * | momOdd | ||
| ) |
Definition at line 629 of file fermion_force_quda.cu.
| void fermion_force_cuda | ( | double | eps, |
| double | weight1, | ||
| double | weight2, | ||
| void * | act_path_coeff, | ||
| FullHw | cudaHw, | ||
| cudaGaugeField & | siteLink, | ||
| cudaGaugeField & | cudaMom, | ||
| QudaGaugeParam * | param | ||
| ) |
Definition at line 1383 of file fermion_force_quda.cu.
| void fermion_force_init_cuda | ( | QudaGaugeParam * | param | ) |
Definition at line 397 of file fermion_force_quda.cu.
1.7.4