QUDA  0.9.0
Macros
io_spinor.h File Reference
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Macros

#define READ_SPINOR_DOUBLE(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_GHOST_DOUBLE(spinor_, stride, sp_idx, norm_idx, dir)
 
#define READ_SPINOR_DOUBLE_UP(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_DOUBLE_DOWN(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_SINGLE(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_GHOST_SINGLE(spinor_, stride, sp_idx, norm_idx, dir)
 
#define READ_SPINOR_SINGLE_UP(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_SINGLE_DOWN(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF_(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF(spinor, stride, sp_idx, norm_idx)   READ_SPINOR_HALF_(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_GHOST_HALF_(spinor_, stride, sp_idx, norm_idx, dir)
 
#define READ_SPINOR_GHOST_HALF(spinor, stride, sp_idx, norm_idx, dir)   READ_SPINOR_GHOST_HALF_(spinor, stride, sp_idx, norm_idx, dir)
 
#define READ_SPINOR_HALF_UP_(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF_UP(spinor, stride, sp_idx, norm_idx)   READ_SPINOR_HALF_UP_(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF_DOWN_(spinor_, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF_DOWN(spinor, stride, sp_idx, norm_idx)   READ_SPINOR_HALF_DOWN_(spinor, stride, sp_idx, norm_idx)
 
#define READ_ACCUM_DOUBLE(spinor_, stride)
 
#define READ_ACCUM_SINGLE(spinor_, stride)
 
#define READ_ACCUM_HALF_(spinor_, stride)
 
#define READ_ACCUM_HALF(spinor, stride)   READ_ACCUM_HALF_(spinor, stride)
 
#define READ_SPINOR_DOUBLE_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_GHOST_DOUBLE_TEX(spinor, stride, sp_idx, norm_idx, dir)
 
#define READ_SPINOR_DOUBLE_UP_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_DOUBLE_DOWN_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_ACCUM_DOUBLE_TEX(spinor, stride)
 
#define READ_SPINOR_SINGLE_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_GHOST_SINGLE_TEX(spinor, stride, sp_idx, norm_idx, dir)
 
#define READ_SPINOR_SINGLE_UP_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_SINGLE_DOWN_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF_TEX_(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF_TEX(spinor, stride, sp_idx, norm_idx)   READ_SPINOR_HALF_TEX_(spinor, stride, sp_idx, norm_idx) \
 
#define READ_SPINOR_GHOST_HALF_TEX_(spinor, stride, sp_idx, norm_idx, dir)
 
#define READ_SPINOR_GHOST_HALF_TEX(spinor, stride, sp_idx, norm_idx, dir)   READ_SPINOR_GHOST_HALF_TEX_(spinor, stride, sp_idx, norm_idx, dir) \
 
#define READ_SPINOR_HALF_UP_TEX_(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF_UP_TEX(spinor, stride, sp_idx, norm_idx)   READ_SPINOR_HALF_UP_TEX_(spinor, stride, sp_idx, norm_idx) \
 
#define READ_SPINOR_HALF_DOWN_TEX_(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_HALF_DOWN_TEX(spinor, stride, sp_idx, norm_idx)   READ_SPINOR_HALF_DOWN_TEX_(spinor, stride, sp_idx, norm_idx) \
 
#define READ_ACCUM_SINGLE_TEX(spinor, stride)
 
#define READ_ACCUM_HALF_TEX_(spinor, stride)
 
#define READ_ACCUM_HALF_TEX(spinor, stride)   READ_ACCUM_HALF_TEX_(spinor, stride)
 
#define WRITE_SPINOR_DOUBLE2(stride)
 
#define WRITE_SPINOR_FLOAT4(stride)
 
#define WRITE_SPINOR_SHORT4(stride)
 
#define WRITE_SPINOR_DOUBLE2_STR(stride)
 
#define WRITE_SPINOR_FLOAT4_STR(stride)
 
#define WRITE_SPINOR_SHORT4_STR(stride)
 
#define READ_HALF_SPINOR   READ_SPINOR_UP
 
#define WRITE_HALF_SPINOR_DOUBLE2(stride, sid)
 
#define WRITE_HALF_SPINOR_FLOAT4(stride, sid)
 
#define WRITE_HALF_SPINOR_SHORT4(stride, sid)
 
#define WRITE_FLAVOR_SPINOR_DOUBLE2()
 ndeg tm: More...
 
#define WRITE_FLAVOR_SPINOR_FLOAT4()
 
#define WRITE_FLAVOR_SPINOR_SHORT4()
 
#define READ_1ST_NBR_SPINOR_DOUBLE_TEX(spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_GHOST_DOUBLE_TEX(spinor, idx, mystride, dir)
 
#define READ_KS_NBR_SPINOR_DOUBLE_TEX(T, spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_GHOST_DOUBLE_TEX(T, spinor, idx, mystride, dir)
 
#define READ_1ST_NBR_SPINOR_SINGLE_TEX(spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_GHOST_SINGLE_TEX(spinor, idx, mystride, dir)
 
#define READ_KS_NBR_SPINOR_SINGLE_TEX(T, spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_GHOST_SINGLE_TEX(T, spinor, idx, mystride, dir)
 
#define READ_1ST_NBR_SPINOR_HALF_TEX_(spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_GHOST_HALF_TEX_(spinor, idx, mystride, dir)
 
#define READ_1ST_NBR_SPINOR_HALF_TEX(spinor, idx, mystride)   READ_1ST_NBR_SPINOR_HALF_TEX_(spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_GHOST_HALF_TEX(spinor, idx, mystride, dir)   READ_1ST_NBR_SPINOR_GHOST_HALF_TEX_(spinor, idx, mystride, dir)
 
#define READ_KS_NBR_SPINOR_HALF_TEX_(T, spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_GHOST_HALF_TEX_(T, spinor, idx, mystride, dir)
 
#define READ_KS_NBR_SPINOR_HALF_TEX(T, spinor, idx, mystride)   READ_KS_NBR_SPINOR_HALF_TEX_(T, spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_GHOST_HALF_TEX(T, spinor, idx, mystride, dir)   READ_KS_NBR_SPINOR_GHOST_HALF_TEX_(T, spinor, idx, mystride, dir)
 
#define READ_1ST_NBR_SPINOR_DOUBLE(spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_GHOST_DOUBLE(spinor, idx, mystride, dir)
 
#define READ_KS_NBR_SPINOR_DOUBLE(T, spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_SINGLE(spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_GHOST_SINGLE(spinor, idx, mystride, dir)
 
#define READ_KS_NBR_SPINOR_SINGLE(T, spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_HALF(spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_HALF(T, spinor, idx, mystride)
 
#define WRITE_ST_SPINOR_DOUBLE2(out, sid, mystride)
 
#define WRITE_ST_SPINOR_FLOAT2(out, sid, mystride)
 
#define WRITE_ST_SPINOR_SHORT2(out, sid, mystride)
 
#define WRITE_ST_SPINOR_DOUBLE2_STR(out, sid, mystride)
 
#define WRITE_ST_SPINOR_FLOAT2_STR(out, sid, mystride)
 
#define WRITE_ST_SPINOR_SHORT2_STR(out, sid, mystride)
 
#define READ_AND_SUM_ST_SPINOR_DOUBLE_TEX(spinor, sid)
 
#define READ_AND_SUM_ST_SPINOR_SINGLE_TEX(spinor, sid)
 
#define READ_AND_SUM_ST_SPINOR_HALF_TEX_(spinor, sid)
 
#define READ_AND_SUM_ST_SPINOR_HALF_TEX(spinor, sid)   READ_AND_SUM_ST_SPINOR_HALF_TEX_(spinor,sid)
 
#define READ_AND_SUM_ST_SPINOR(spinor, sid)
 
#define READ_AND_SUM_ST_SPINOR_HALF_(spinor, sid)
 
#define READ_AND_SUM_ST_SPINOR_HALF(spinor, sid)   READ_AND_SUM_ST_SPINOR_HALF_(spinor,sid)
 
#define READ_ST_ACCUM_DOUBLE_TEX(spinor, sid)
 
#define READ_ST_ACCUM_SINGLE_TEX(spinor, sid)
 
#define READ_ST_ACCUM_HALF_TEX_(spinor, sid)
 
#define READ_ST_ACCUM_HALF_TEX(spinor, sid)   READ_ST_ACCUM_HALF_TEX_(spinor,sid)
 
#define READ_ST_ACCUM_DOUBLE(spinor, sid)
 
#define READ_ST_ACCUM_SINGLE(spinor, sid)
 
#define READ_ST_ACCUM_HALF(spinor, sid)
 
#define WRITE_SPINOR_SHARED_REAL(tx, ty, tz, reg)
 
#define WRITE_SPINOR_SHARED_DOUBLE2   WRITE_SPINOR_SHARED_REAL
 
#define READ_SPINOR_SHARED_DOUBLE2(tx, ty, tz)
 
#define WRITE_SPINOR_SHARED_FLOAT4   WRITE_SPINOR_SHARED_REAL
 
#define READ_SPINOR_SHARED_FLOAT4(tx, ty, tz)
 
#define READ_ACCUM_FLAVOR_DOUBLE(spinor, stride, fl_stride)
 **************************only for ndeg tm:****************************** More...
 
#define READ_ACCUM_FLAVOR_SINGLE(spinor, stride, flv_stride)
 
#define READ_ACCUM_FLAVOR_HALF_(spinor, stride, flv_stride)
 
#define READ_ACCUM_FLAVOR_HALF(spinor, stride, flv_stride)   READ_ACCUM_FLAVOR_HALF_(spinor, stride, flv_stride)
 
#define READ_ACCUM_FLAVOR_DOUBLE_TEX(spinor, stride, flv_stride)
 
#define READ_ACCUM_FLAVOR_SINGLE_TEX(spinor, stride, flv_stride)
 
#define READ_ACCUM_HALF_FLAVOR_TEX_(spinor, stride, flv_stride)
 
#define READ_ACCUM_FLAVOR_HALF_TEX(spinor, stride, flv_stride)   READ_ACCUM_HALF_FLAVOR_TEX_(spinor, stride, flv_stride)
 
#define ASSN_ACCUM_DOUBLE(spinor, stride, fl_stride)
 
#define ASSN_ACCUM_SINGLE(spinor, stride, fl_stride)
 
#define ASSN_ACCUM_HALF_(spinor, stride, fl_stride)
 
#define ASSN_ACCUM_HALF(spinor, stride, fl_stride)   ASSN_ACCUM_HALF_(spinor, stride, fl_stride)
 
#define ASSN_ACCUM_DOUBLE_TEX(spinor, stride, fl_stride)
 
#define ASSN_ACCUM_SINGLE_TEX(spinor, stride, fl_stride)
 
#define ASSN_ACCUM_HALF_TEX_(spinor, stride, fl_stride)
 
#define ASSN_ACCUM_HALF_TEX(spinor, stride, fl_stride)   ASSN_ACCUM_HALF_TEX_(spinor, stride, fl_stride)
 
#define APPLY_TWIST_INV(a, b, reg)
 **************************only for deg tm:******************************* More...
 
#define APPLY_TWIST(a, reg)
 

Macro Definition Documentation

◆ APPLY_TWIST

#define APPLY_TWIST (   a,
  reg 
)

Definition at line 1379 of file io_spinor.h.

◆ APPLY_TWIST_INV

#define APPLY_TWIST_INV (   a,
  b,
  reg 
)

**************************only for deg tm:*******************************

Definition at line 1314 of file io_spinor.h.

Referenced by if().

◆ ASSN_ACCUM_DOUBLE

#define ASSN_ACCUM_DOUBLE (   spinor,
  stride,
  fl_stride 
)
Value:
accum0 = spinor[sid + fl_stride + 0*stride]; \
accum1 = spinor[sid + fl_stride + 1*stride]; \
accum2 = spinor[sid + fl_stride + 2*stride]; \
accum3 = spinor[sid + fl_stride + 3*stride]; \
accum4 = spinor[sid + fl_stride + 4*stride]; \
accum5 = spinor[sid + fl_stride + 5*stride]; \
accum6 = spinor[sid + fl_stride + 6*stride]; \
accum7 = spinor[sid + fl_stride + 7*stride]; \
accum8 = spinor[sid + fl_stride + 8*stride]; \
accum9 = spinor[sid + fl_stride + 9*stride]; \
accum10 = spinor[sid + fl_stride + 10*stride]; \
accum11 = spinor[sid + fl_stride + 11*stride];
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 1225 of file io_spinor.h.

◆ ASSN_ACCUM_DOUBLE_TEX

#define ASSN_ACCUM_DOUBLE_TEX (   spinor,
  stride,
  fl_stride 
)
Value:
accum0 = fetch_double2((spinor), sid + fl_stride + 0*(stride)); \
accum1 = fetch_double2((spinor), sid + fl_stride + 1*(stride)); \
accum2 = fetch_double2((spinor), sid + fl_stride + 2*(stride)); \
accum3 = fetch_double2((spinor), sid + fl_stride + 3*(stride)); \
accum4 = fetch_double2((spinor), sid + fl_stride + 4*(stride)); \
accum5 = fetch_double2((spinor), sid + fl_stride + 5*(stride)); \
accum6 = fetch_double2((spinor), sid + fl_stride + 6*(stride)); \
accum7 = fetch_double2((spinor), sid + fl_stride + 7*(stride)); \
accum8 = fetch_double2((spinor), sid + fl_stride + 8*(stride)); \
accum9 = fetch_double2((spinor), sid + fl_stride + 9*(stride)); \
accum10 = fetch_double2((spinor), sid + fl_stride + 10*(stride)); \
accum11 = fetch_double2((spinor), sid + fl_stride + 11*(stride));
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 1269 of file io_spinor.h.

◆ ASSN_ACCUM_HALF

#define ASSN_ACCUM_HALF (   spinor,
  stride,
  fl_stride 
)    ASSN_ACCUM_HALF_(spinor, stride, fl_stride)

Definition at line 1265 of file io_spinor.h.

◆ ASSN_ACCUM_HALF_

#define ASSN_ACCUM_HALF_ (   spinor,
  stride,
  fl_stride 
)
Value:
accum0 = short42float4(spinor[sid + fl_stride + 0*stride]); \
accum1 = short42float4(spinor[sid + fl_stride + 1*stride]); \
accum2 = short42float4(spinor[sid + fl_stride + 2*stride]); \
accum3 = short42float4(spinor[sid + fl_stride + 3*stride]); \
accum4 = short42float4(spinor[sid + fl_stride + 4*stride]); \
accum5 = short42float4(spinor[sid + fl_stride + 5*stride]); \
{\
float C = (spinor ## Norm)[sid + fl_stride]; \
accum0.x *= C; accum0.y *= C; accum0.z *= C; accum0.w *= C; \
accum1.x *= C; accum1.y *= C; accum1.z *= C; accum1.w *= C; \
accum2.x *= C; accum2.y *= C; accum2.z *= C; accum2.w *= C; \
accum3.x *= C; accum3.y *= C; accum3.z *= C; accum3.w *= C; \
accum4.x *= C; accum4.y *= C; accum4.z *= C; accum4.w *= C; \
accum5.x *= C; accum5.y *= C; accum5.z *= C; accum5.w *= C; \
}
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 1247 of file io_spinor.h.

◆ ASSN_ACCUM_HALF_TEX

#define ASSN_ACCUM_HALF_TEX (   spinor,
  stride,
  fl_stride 
)    ASSN_ACCUM_HALF_TEX_(spinor, stride, fl_stride)

Definition at line 1309 of file io_spinor.h.

◆ ASSN_ACCUM_HALF_TEX_

#define ASSN_ACCUM_HALF_TEX_ (   spinor,
  stride,
  fl_stride 
)
Value:
accum0 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 0*(stride)); \
accum1 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 1*(stride)); \
accum2 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 2*(stride)); \
accum3 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 3*(stride)); \
accum4 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 4*(stride)); \
accum5 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 5*(stride)); \
{\
float C = TEX1DFETCH(float, (spinor ## Norm), sid + fl_stride); \
accum0.x *= C; accum0.y *= C; accum0.z *= C; accum0.w *= C; \
accum1.x *= C; accum1.y *= C; accum1.z *= C; accum1.w *= C; \
accum2.x *= C; accum2.y *= C; accum2.z *= C; accum2.w *= C; \
accum3.x *= C; accum3.y *= C; accum3.z *= C; accum3.w *= C; \
accum4.x *= C; accum4.y *= C; accum4.z *= C; accum4.w *= C; \
accum5.x *= C; accum5.y *= C; accum5.z *= C; accum5.w *= C; \
}
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 1292 of file io_spinor.h.

◆ ASSN_ACCUM_SINGLE

#define ASSN_ACCUM_SINGLE (   spinor,
  stride,
  fl_stride 
)
Value:
accum0 = spinor[sid + fl_stride + 0*(stride)]; \
accum1 = spinor[sid + fl_stride + 1*(stride)]; \
accum2 = spinor[sid + fl_stride + 2*(stride)]; \
accum3 = spinor[sid + fl_stride + 3*(stride)]; \
accum4 = spinor[sid + fl_stride + 4*(stride)]; \
accum5 = spinor[sid + fl_stride + 5*(stride)];
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 1239 of file io_spinor.h.

◆ ASSN_ACCUM_SINGLE_TEX

#define ASSN_ACCUM_SINGLE_TEX (   spinor,
  stride,
  fl_stride 
)
Value:
accum0 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 0*(stride)); \
accum1 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 1*(stride)); \
accum2 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 2*(stride)); \
accum3 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 3*(stride)); \
accum4 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 4*(stride)); \
accum5 = TEX1DFETCH(float4, (spinor), sid + fl_stride + 5*(stride));
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 1284 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_DOUBLE

#define READ_1ST_NBR_SPINOR_DOUBLE (   spinor,
  idx,
  mystride 
)
Value:
double2 I0 = spinor[idx + 0*mystride]; \
double2 I1 = spinor[idx + 1*mystride]; \
double2 I2 = spinor[idx + 2*mystride];
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 789 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_DOUBLE_TEX

#define READ_1ST_NBR_SPINOR_DOUBLE_TEX (   spinor,
  idx,
  mystride 
)
Value:
double2 I0 = fetch_double2((spinor), idx + 0*mystride); \
double2 I1 = fetch_double2((spinor), idx + 1*mystride); \
double2 I2 = fetch_double2((spinor), idx + 2*mystride);
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 645 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_GHOST_DOUBLE

#define READ_1ST_NBR_SPINOR_GHOST_DOUBLE (   spinor,
  idx,
  mystride,
  dir 
)
Value:
double2 I0 = spinor[dir][idx + 0*mystride]; \
double2 I1 = spinor[dir][idx + 1*mystride]; \
double2 I2 = spinor[dir][idx + 2*mystride];
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 794 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_GHOST_DOUBLE_TEX

#define READ_1ST_NBR_SPINOR_GHOST_DOUBLE_TEX (   spinor,
  idx,
  mystride,
  dir 
)
Value:
double2 I0 = fetch_double2((spinor), idx + 0*mystride); \
double2 I1 = fetch_double2((spinor), idx + 1*mystride); \
double2 I2 = fetch_double2((spinor), idx + 2*mystride);
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 656 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_GHOST_HALF_TEX

#define READ_1ST_NBR_SPINOR_GHOST_HALF_TEX (   spinor,
  idx,
  mystride,
  dir 
)    READ_1ST_NBR_SPINOR_GHOST_HALF_TEX_(spinor, idx, mystride, dir)

Definition at line 748 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_GHOST_HALF_TEX_

#define READ_1ST_NBR_SPINOR_GHOST_HALF_TEX_ (   spinor,
  idx,
  mystride,
  dir 
)
Value:
float2 I0 = TEX1DFETCH(float2, (spinor), idx + 0*mystride); \
float2 I1 = TEX1DFETCH(float2, (spinor), idx + 1*mystride); \
float2 I2 = TEX1DFETCH(float2, (spinor), idx + 2*mystride); \
{ \
float C = TEX1DFETCH(float, (spinor ## Norm), norm_idx1); \
I0.x *= C; I0.y *= C; \
I1.x *= C; I1.y *= C; \
I2.x *= C; I2.y *= C;}
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 734 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_GHOST_SINGLE

#define READ_1ST_NBR_SPINOR_GHOST_SINGLE (   spinor,
  idx,
  mystride,
  dir 
)
Value:
float2 I0 = spinor[dir][idx + 0*mystride]; \
float2 I1 = spinor[dir][idx + 1*mystride]; \
float2 I2 = spinor[dir][idx + 2*mystride];
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 809 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_GHOST_SINGLE_TEX

#define READ_1ST_NBR_SPINOR_GHOST_SINGLE_TEX (   spinor,
  idx,
  mystride,
  dir 
)
Value:
float2 I0 = TEX1DFETCH(float2, (spinor), idx + 0*mystride); \
float2 I1 = TEX1DFETCH(float2, (spinor), idx + 1*mystride); \
float2 I2 = TEX1DFETCH(float2, (spinor), idx + 2*mystride);
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 690 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_HALF

#define READ_1ST_NBR_SPINOR_HALF (   spinor,
  idx,
  mystride 
)
Value:
float2 I0, I1, I2; \
{ \
short2 S0 = in[idx + 0*mystride]; \
short2 S1 = in[idx + 1*mystride]; \
short2 S2 = in[idx + 2*mystride]; \
float C = inNorm[idx]; \
I0.x =C*short2float(S0.x); I0.y =C*short2float(S0.y); \
I1.x =C*short2float(S1.x); I1.y =C*short2float(S1.y); \
I2.x =C*short2float(S2.x); I2.y =C*short2float(S2.y); \
}
in
cpuColorSpinorField * in
Definition: staggered_invert_test.cpp:44
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355

Definition at line 819 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_HALF_TEX

#define READ_1ST_NBR_SPINOR_HALF_TEX (   spinor,
  idx,
  mystride 
)    READ_1ST_NBR_SPINOR_HALF_TEX_(spinor, idx, mystride)

Definition at line 745 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_HALF_TEX_

#define READ_1ST_NBR_SPINOR_HALF_TEX_ (   spinor,
  idx,
  mystride 
)
Value:
float2 I0 = TEX1DFETCH(float2, (spinor), idx + 0*mystride); \
float2 I1 = TEX1DFETCH(float2, (spinor), idx + 1*mystride); \
float2 I2 = TEX1DFETCH(float2, (spinor), idx + 2*mystride); \
{ \
float C = TEX1DFETCH(float, (spinor ## Norm), norm_idx1); \
I0.x *= C; I0.y *= C; \
I1.x *= C; I1.y *= C; \
I2.x *= C; I2.y *= C;}
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 713 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_SINGLE

#define READ_1ST_NBR_SPINOR_SINGLE (   spinor,
  idx,
  mystride 
)
Value:
float2 I0 = spinor[idx + 0*mystride]; \
float2 I1 = spinor[idx + 1*mystride]; \
float2 I2 = spinor[idx + 2*mystride];
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 804 of file io_spinor.h.

◆ READ_1ST_NBR_SPINOR_SINGLE_TEX

#define READ_1ST_NBR_SPINOR_SINGLE_TEX (   spinor,
  idx,
  mystride 
)
Value:
float2 I0 = TEX1DFETCH(float2, (spinor), idx + 0*mystride); \
float2 I1 = TEX1DFETCH(float2, (spinor), idx + 1*mystride); \
float2 I2 = TEX1DFETCH(float2, (spinor), idx + 2*mystride);
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 679 of file io_spinor.h.

◆ READ_ACCUM_DOUBLE

#define READ_ACCUM_DOUBLE (   spinor_,
  stride 
)
Value:
double2 *spinor = (double2*)spinor_; \
double2 accum0 = spinor[sid + 0*stride]; \
double2 accum1 = spinor[sid + 1*stride]; \
double2 accum2 = spinor[sid + 2*stride]; \
double2 accum3 = spinor[sid + 3*stride]; \
double2 accum4 = spinor[sid + 4*stride]; \
double2 accum5 = spinor[sid + 5*stride]; \
double2 accum6 = spinor[sid + 6*stride]; \
double2 accum7 = spinor[sid + 7*stride]; \
double2 accum8 = spinor[sid + 8*stride]; \
double2 accum9 = spinor[sid + 9*stride]; \
double2 accum10 = spinor[sid + 10*stride]; \
double2 accum11 = spinor[sid + 11*stride];
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 128 of file io_spinor.h.

◆ READ_ACCUM_DOUBLE_TEX

#define READ_ACCUM_DOUBLE_TEX (   spinor,
  stride 
)
Value:
double2 accum0 = fetch_double2((spinor), sid + 0*(stride)); \
double2 accum1 = fetch_double2((spinor), sid + 1*(stride)); \
double2 accum2 = fetch_double2((spinor), sid + 2*(stride)); \
double2 accum3 = fetch_double2((spinor), sid + 3*(stride)); \
double2 accum4 = fetch_double2((spinor), sid + 4*(stride)); \
double2 accum5 = fetch_double2((spinor), sid + 5*(stride)); \
double2 accum6 = fetch_double2((spinor), sid + 6*(stride)); \
double2 accum7 = fetch_double2((spinor), sid + 7*(stride)); \
double2 accum8 = fetch_double2((spinor), sid + 8*(stride)); \
double2 accum9 = fetch_double2((spinor), sid + 9*(stride)); \
double2 accum10 = fetch_double2((spinor), sid + 10*(stride)); \
double2 accum11 = fetch_double2((spinor), sid + 11*(stride));
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 218 of file io_spinor.h.

◆ READ_ACCUM_FLAVOR_DOUBLE

#define READ_ACCUM_FLAVOR_DOUBLE (   spinor,
  stride,
  fl_stride 
)
Value:
double2 flv1_accum0 = spinor[sid + 0*stride]; \
double2 flv1_accum1 = spinor[sid + 1*stride]; \
double2 flv1_accum2 = spinor[sid + 2*stride]; \
double2 flv1_accum3 = spinor[sid + 3*stride]; \
double2 flv1_accum4 = spinor[sid + 4*stride]; \
double2 flv1_accum5 = spinor[sid + 5*stride]; \
double2 flv1_accum6 = spinor[sid + 6*stride]; \
double2 flv1_accum7 = spinor[sid + 7*stride]; \
double2 flv1_accum8 = spinor[sid + 8*stride]; \
double2 flv1_accum9 = spinor[sid + 9*stride]; \
double2 flv1_accum10 = spinor[sid + 10*stride]; \
double2 flv1_accum11 = spinor[sid + 11*stride]; \
double2 flv2_accum0 = spinor[sid + fl_stride + 0*stride]; \
double2 flv2_accum1 = spinor[sid + fl_stride + 1*stride]; \
double2 flv2_accum2 = spinor[sid + fl_stride + 2*stride]; \
double2 flv2_accum3 = spinor[sid + fl_stride + 3*stride]; \
double2 flv2_accum4 = spinor[sid + fl_stride + 4*stride]; \
double2 flv2_accum5 = spinor[sid + fl_stride + 5*stride]; \
double2 flv2_accum6 = spinor[sid + fl_stride + 6*stride]; \
double2 flv2_accum7 = spinor[sid + fl_stride + 7*stride]; \
double2 flv2_accum8 = spinor[sid + fl_stride + 8*stride]; \
double2 flv2_accum9 = spinor[sid + fl_stride + 9*stride]; \
double2 flv2_accum10 = spinor[sid + fl_stride + 10*stride]; \
double2 flv2_accum11 = spinor[sid + fl_stride + 11*stride];
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

**************************only for ndeg tm:******************************

Definition at line 1078 of file io_spinor.h.

◆ READ_ACCUM_FLAVOR_DOUBLE_TEX

#define READ_ACCUM_FLAVOR_DOUBLE_TEX (   spinor,
  stride,
  flv_stride 
)
Value:
double2 flv1_accum0 = fetch_double2((spinor), sid + 0*(stride)); \
double2 flv1_accum1 = fetch_double2((spinor), sid + 1*(stride)); \
double2 flv1_accum2 = fetch_double2((spinor), sid + 2*(stride)); \
double2 flv1_accum3 = fetch_double2((spinor), sid + 3*(stride)); \
double2 flv1_accum4 = fetch_double2((spinor), sid + 4*(stride)); \
double2 flv1_accum5 = fetch_double2((spinor), sid + 5*(stride)); \
double2 flv1_accum6 = fetch_double2((spinor), sid + 6*(stride)); \
double2 flv1_accum7 = fetch_double2((spinor), sid + 7*(stride)); \
double2 flv1_accum8 = fetch_double2((spinor), sid + 8*(stride)); \
double2 flv1_accum9 = fetch_double2((spinor), sid + 9*(stride)); \
double2 flv1_accum10 = fetch_double2((spinor), sid + 10*(stride)); \
double2 flv1_accum11 = fetch_double2((spinor), sid + 11*(stride)); \
double2 flv2_accum0 = fetch_double2((spinor), sid + flv_stride + 0*(stride)); \
double2 flv2_accum1 = fetch_double2((spinor), sid + flv_stride + 1*(stride)); \
double2 flv2_accum2 = fetch_double2((spinor), sid + flv_stride + 2*(stride)); \
double2 flv2_accum3 = fetch_double2((spinor), sid + flv_stride + 3*(stride)); \
double2 flv2_accum4 = fetch_double2((spinor), sid + flv_stride + 4*(stride)); \
double2 flv2_accum5 = fetch_double2((spinor), sid + flv_stride + 5*(stride)); \
double2 flv2_accum6 = fetch_double2((spinor), sid + flv_stride + 6*(stride)); \
double2 flv2_accum7 = fetch_double2((spinor), sid + flv_stride + 7*(stride)); \
double2 flv2_accum8 = fetch_double2((spinor), sid + flv_stride + 8*(stride)); \
double2 flv2_accum9 = fetch_double2((spinor), sid + flv_stride + 9*(stride)); \
double2 flv2_accum10 = fetch_double2((spinor), sid + flv_stride + 10*(stride)); \
double2 flv2_accum11 = fetch_double2((spinor), sid + flv_stride + 11*(stride));
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 1151 of file io_spinor.h.

◆ READ_ACCUM_FLAVOR_HALF

#define READ_ACCUM_FLAVOR_HALF (   spinor,
  stride,
  flv_stride 
)    READ_ACCUM_FLAVOR_HALF_(spinor, stride, flv_stride)

Definition at line 1148 of file io_spinor.h.

◆ READ_ACCUM_FLAVOR_HALF_

#define READ_ACCUM_FLAVOR_HALF_ (   spinor,
  stride,
  flv_stride 
)
Value:
float4 flv1_accum0 = short42float4(spinor[sid + 0*stride]); \
float4 flv1_accum1 = short42float4(spinor[sid + 1*stride]); \
float4 flv1_accum2 = short42float4(spinor[sid + 2*stride]); \
float4 flv1_accum3 = short42float4(spinor[sid + 3*stride]); \
float4 flv1_accum4 = short42float4(spinor[sid + 4*stride]); \
float4 flv1_accum5 = short42float4(spinor[sid + 5*stride]); \
float C = (spinor ## Norm)[sid]; \
flv1_accum0.x *= C; flv1_accum0.y *= C; flv1_accum0.z *= C; flv1_accum0.w *= C; \
flv1_accum1.x *= C; flv1_accum1.y *= C; flv1_accum1.z *= C; flv1_accum1.w *= C; \
flv1_accum2.x *= C; flv1_accum2.y *= C; flv1_accum2.z *= C; flv1_accum2.w *= C; \
flv1_accum3.x *= C; flv1_accum3.y *= C; flv1_accum3.z *= C; flv1_accum3.w *= C; \
flv1_accum4.x *= C; flv1_accum4.y *= C; flv1_accum4.z *= C; flv1_accum4.w *= C; \
flv1_accum5.x *= C; flv1_accum5.y *= C; flv1_accum5.z *= C; flv1_accum5.w *= C; \
float4 flv2_accum0 = short42float4(spinor[sid + flv_stride + 0*stride]); \
float4 flv2_accum1 = short42float4(spinor[sid + flv_stride + 1*stride]); \
float4 flv2_accum2 = short42float4(spinor[sid + flv_stride + 2*stride]); \
float4 flv2_accum3 = short42float4(spinor[sid + flv_stride + 3*stride]); \
float4 flv2_accum4 = short42float4(spinor[sid + flv_stride + 4*stride]); \
float4 flv2_accum5 = short42float4(spinor[sid + flv_stride + 5*stride]); \
C = (spinor ## Norm)[sid + fl_stride]; \
flv2_accum0.x *= C; flv2_accum0.y *= C; flv2_accum0.z *= C; flv2_accum0.w *= C; \
flv2_accum1.x *= C; flv2_accum1.y *= C; flv2_accum1.z *= C; flv2_accum1.w *= C; \
flv2_accum2.x *= C; flv2_accum2.y *= C; flv2_accum2.z *= C; flv2_accum2.w *= C; \
flv2_accum3.x *= C; flv2_accum3.y *= C; flv2_accum3.z *= C; flv2_accum3.w *= C; \
flv2_accum4.x *= C; flv2_accum4.y *= C; flv2_accum4.z *= C; flv2_accum4.w *= C; \
flv2_accum5.x *= C; flv2_accum5.y *= C; flv2_accum5.z *= C; flv2_accum5.w *= C;
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 1120 of file io_spinor.h.

◆ READ_ACCUM_FLAVOR_HALF_TEX

#define READ_ACCUM_FLAVOR_HALF_TEX (   spinor,
  stride,
  flv_stride 
)    READ_ACCUM_HALF_FLAVOR_TEX_(spinor, stride, flv_stride)

Definition at line 1221 of file io_spinor.h.

◆ READ_ACCUM_FLAVOR_SINGLE

#define READ_ACCUM_FLAVOR_SINGLE (   spinor,
  stride,
  flv_stride 
)
Value:
float4 flv1_accum0 = spinor[sid + 0*(stride)]; \
float4 flv1_accum1 = spinor[sid + 1*(stride)]; \
float4 flv1_accum2 = spinor[sid + 2*(stride)]; \
float4 flv1_accum3 = spinor[sid + 3*(stride)]; \
float4 flv1_accum4 = spinor[sid + 4*(stride)]; \
float4 flv1_accum5 = spinor[sid + 5*(stride)]; \
float4 flv2_accum0 = spinor[sid + flv_stride + 0*(stride)]; \
float4 flv2_accum1 = spinor[sid + flv_stride + 1*(stride)]; \
float4 flv2_accum2 = spinor[sid + flv_stride + 2*(stride)]; \
float4 flv2_accum3 = spinor[sid + flv_stride + 3*(stride)]; \
float4 flv2_accum4 = spinor[sid + flv_stride + 4*(stride)]; \
float4 flv2_accum5 = spinor[sid + flv_stride + 5*(stride)];
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 1105 of file io_spinor.h.

◆ READ_ACCUM_FLAVOR_SINGLE_TEX

#define READ_ACCUM_FLAVOR_SINGLE_TEX (   spinor,
  stride,
  flv_stride 
)
Value:
float4 flv1_accum0 = TEX1DFETCH(float4, (spinor), sid + 0*(stride)); \
float4 flv1_accum1 = TEX1DFETCH(float4, (spinor), sid + 1*(stride)); \
float4 flv1_accum2 = TEX1DFETCH(float4, (spinor), sid + 2*(stride)); \
float4 flv1_accum3 = TEX1DFETCH(float4, (spinor), sid + 3*(stride)); \
float4 flv1_accum4 = TEX1DFETCH(float4, (spinor), sid + 4*(stride)); \
float4 flv1_accum5 = TEX1DFETCH(float4, (spinor), sid + 5*(stride)); \
float4 flv2_accum0 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 0*(stride)); \
float4 flv2_accum1 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 1*(stride)); \
float4 flv2_accum2 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 2*(stride)); \
float4 flv2_accum3 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 3*(stride)); \
float4 flv2_accum4 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 4*(stride)); \
float4 flv2_accum5 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 5*(stride));
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 1178 of file io_spinor.h.

◆ READ_ACCUM_HALF

#define READ_ACCUM_HALF (   spinor,
  stride 
)    READ_ACCUM_HALF_(spinor, stride)

Definition at line 168 of file io_spinor.h.

◆ READ_ACCUM_HALF_

#define READ_ACCUM_HALF_ (   spinor_,
  stride 
)
Value:
short4 *spinor = (short4*)spinor_; \
float4 accum0 = short42float4(spinor[sid + 0*stride]); \
float4 accum1 = short42float4(spinor[sid + 1*stride]); \
float4 accum2 = short42float4(spinor[sid + 2*stride]); \
float4 accum3 = short42float4(spinor[sid + 3*stride]); \
float4 accum4 = short42float4(spinor[sid + 4*stride]); \
float4 accum5 = short42float4(spinor[sid + 5*stride]); \
float C = (spinor_ ## Norm)[sid]; \
accum0.x *= C; accum0.y *= C; accum0.z *= C; accum0.w *= C; \
accum1.x *= C; accum1.y *= C; accum1.z *= C; accum1.w *= C; \
accum2.x *= C; accum2.y *= C; accum2.z *= C; accum2.w *= C; \
accum3.x *= C; accum3.y *= C; accum3.z *= C; accum3.w *= C; \
accum4.x *= C; accum4.y *= C; accum4.z *= C; accum4.w *= C; \
accum5.x *= C; accum5.y *= C; accum5.z *= C; accum5.w *= C;
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 152 of file io_spinor.h.

◆ READ_ACCUM_HALF_FLAVOR_TEX_

#define READ_ACCUM_HALF_FLAVOR_TEX_ (   spinor,
  stride,
  flv_stride 
)
Value:
float4 flv1_accum0 = TEX1DFETCH(float4, (spinor), sid + 0*(stride)); \
float4 flv1_accum1 = TEX1DFETCH(float4, (spinor), sid + 1*(stride)); \
float4 flv1_accum2 = TEX1DFETCH(float4, (spinor), sid + 2*(stride)); \
float4 flv1_accum3 = TEX1DFETCH(float4, (spinor), sid + 3*(stride)); \
float4 flv1_accum4 = TEX1DFETCH(float4, (spinor), sid + 4*(stride)); \
float4 flv1_accum5 = TEX1DFETCH(float4, (spinor), sid + 5*(stride)); \
float C = TEX1DFETCH(float, (spinor ## Norm), sid); \
flv1_accum0.x *= C; flv1_accum0.y *= C; flv1_accum0.z *= C; flv1_accum0.w *= C; \
flv1_accum1.x *= C; flv1_accum1.y *= C; flv1_accum1.z *= C; flv1_accum1.w *= C; \
flv1_accum2.x *= C; flv1_accum2.y *= C; flv1_accum2.z *= C; flv1_accum2.w *= C; \
flv1_accum3.x *= C; flv1_accum3.y *= C; flv1_accum3.z *= C; flv1_accum3.w *= C; \
flv1_accum4.x *= C; flv1_accum4.y *= C; flv1_accum4.z *= C; flv1_accum4.w *= C; \
flv1_accum5.x *= C; flv1_accum5.y *= C; flv1_accum5.z *= C; flv1_accum5.w *= C; \
float4 flv2_accum0 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 0*(stride)); \
float4 flv2_accum1 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 1*(stride)); \
float4 flv2_accum2 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 2*(stride)); \
float4 flv2_accum3 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 3*(stride)); \
float4 flv2_accum4 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 4*(stride)); \
float4 flv2_accum5 = TEX1DFETCH(float4, (spinor), sid + flv_stride + 5*(stride)); \
C = TEX1DFETCH(float, (spinor ## Norm), sid + flv_stride); \
flv2_accum0.x *= C; flv2_accum0.y *= C; flv2_accum0.z *= C; flv2_accum0.w *= C; \
flv2_accum1.x *= C; flv2_accum1.y *= C; flv2_accum1.z *= C; flv2_accum1.w *= C; \
flv2_accum2.x *= C; flv2_accum2.y *= C; flv2_accum2.z *= C; flv2_accum2.w *= C; \
flv2_accum3.x *= C; flv2_accum3.y *= C; flv2_accum3.z *= C; flv2_accum3.w *= C; \
flv2_accum4.x *= C; flv2_accum4.y *= C; flv2_accum4.z *= C; flv2_accum4.w *= C; \
flv2_accum5.x *= C; flv2_accum5.y *= C; flv2_accum5.z *= C; flv2_accum5.w *= C;
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 1192 of file io_spinor.h.

◆ READ_ACCUM_HALF_TEX

#define READ_ACCUM_HALF_TEX (   spinor,
  stride 
)    READ_ACCUM_HALF_TEX_(spinor, stride)

Definition at line 350 of file io_spinor.h.

◆ READ_ACCUM_HALF_TEX_

#define READ_ACCUM_HALF_TEX_ (   spinor,
  stride 
)
Value:
float4 accum0 = TEX1DFETCH(float4, (spinor), sid + 0*(stride)); \
float4 accum1 = TEX1DFETCH(float4, (spinor), sid + 1*(stride)); \
float4 accum2 = TEX1DFETCH(float4, (spinor), sid + 2*(stride)); \
float4 accum3 = TEX1DFETCH(float4, (spinor), sid + 3*(stride)); \
float4 accum4 = TEX1DFETCH(float4, (spinor), sid + 4*(stride)); \
float4 accum5 = TEX1DFETCH(float4, (spinor), sid + 5*(stride)); \
float C = TEX1DFETCH(float, (spinor ## Norm), sid); \
accum0.x *= C; accum0.y *= C; accum0.z *= C; accum0.w *= C; \
accum1.x *= C; accum1.y *= C; accum1.z *= C; accum1.w *= C; \
accum2.x *= C; accum2.y *= C; accum2.z *= C; accum2.w *= C; \
accum3.x *= C; accum3.y *= C; accum3.z *= C; accum3.w *= C; \
accum4.x *= C; accum4.y *= C; accum4.z *= C; accum4.w *= C; \
accum5.x *= C; accum5.y *= C; accum5.z *= C; accum5.w *= C;
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 335 of file io_spinor.h.

◆ READ_ACCUM_SINGLE

#define READ_ACCUM_SINGLE (   spinor_,
  stride 
)
Value:
float4 *spinor = (float4*)spinor_; \
float4 accum0 = spinor[sid + 0*(stride)]; \
float4 accum1 = spinor[sid + 1*(stride)]; \
float4 accum2 = spinor[sid + 2*(stride)]; \
float4 accum3 = spinor[sid + 3*(stride)]; \
float4 accum4 = spinor[sid + 4*(stride)]; \
float4 accum5 = spinor[sid + 5*(stride)];
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 143 of file io_spinor.h.

◆ READ_ACCUM_SINGLE_TEX

#define READ_ACCUM_SINGLE_TEX (   spinor,
  stride 
)
Value:
float4 accum0 = TEX1DFETCH(float4, (spinor), sid + 0*(stride)); \
float4 accum1 = TEX1DFETCH(float4, (spinor), sid + 1*(stride)); \
float4 accum2 = TEX1DFETCH(float4, (spinor), sid + 2*(stride)); \
float4 accum3 = TEX1DFETCH(float4, (spinor), sid + 3*(stride)); \
float4 accum4 = TEX1DFETCH(float4, (spinor), sid + 4*(stride)); \
float4 accum5 = TEX1DFETCH(float4, (spinor), sid + 5*(stride));
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 327 of file io_spinor.h.

◆ READ_AND_SUM_ST_SPINOR

#define READ_AND_SUM_ST_SPINOR (   spinor,
  sid 
)
Value:
o00_re += spinor[0*param.sp_stride+sid].x; o00_im += spinor[0*param.sp_stride+sid].y; \
o01_re += spinor[1*param.sp_stride+sid].x; o01_im += spinor[1*param.sp_stride+sid].y; \
o02_re += spinor[2*param.sp_stride+sid].x; o02_im += spinor[2*param.sp_stride+sid].y; \
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 920 of file io_spinor.h.

◆ READ_AND_SUM_ST_SPINOR_DOUBLE_TEX

#define READ_AND_SUM_ST_SPINOR_DOUBLE_TEX (   spinor,
  sid 
)
Value:
{ \
double2 tmp0 = fetch_double2((spinor), sid + 0*(param.sp_stride)); \
double2 tmp1 = fetch_double2((spinor), sid + 1*(param.sp_stride)); \
double2 tmp2 = fetch_double2((spinor), sid + 2*(param.sp_stride)); \
o00_re += tmp0.x; o00_im += tmp0.y; \
o01_re += tmp1.x; o01_im += tmp1.y; \
o02_re += tmp2.x; o02_im += tmp2.y; }
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
tmp2
#define tmp2
Definition: tmc_core.h:16
tmp1
#define tmp1
Definition: tmc_core.h:15
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359
tmp0
#define tmp0
Definition: tmc_core.h:14
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 892 of file io_spinor.h.

◆ READ_AND_SUM_ST_SPINOR_HALF

#define READ_AND_SUM_ST_SPINOR_HALF (   spinor,
  sid 
)    READ_AND_SUM_ST_SPINOR_HALF_(spinor,sid)

Definition at line 934 of file io_spinor.h.

◆ READ_AND_SUM_ST_SPINOR_HALF_

#define READ_AND_SUM_ST_SPINOR_HALF_ (   spinor,
  sid 
)
Value:
float C = spinor ## Norm[sid]; \
o00_re += C*short2float(spinor[0*param.sp_stride + sid].x); \
o00_im += C*short2float(spinor[0*param.sp_stride + sid].y); \
o01_re += C*short2float(spinor[1*param.sp_stride + sid].x); \
o01_im += C*short2float(spinor[1*param.sp_stride + sid].y); \
o02_re += C*short2float(spinor[2*param.sp_stride + sid].x); \
o02_im += C*short2float(spinor[2*param.sp_stride + sid].y);
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
param
QudaGaugeParam param
Definition: pack_test.cpp:17
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 925 of file io_spinor.h.

◆ READ_AND_SUM_ST_SPINOR_HALF_TEX

#define READ_AND_SUM_ST_SPINOR_HALF_TEX (   spinor,
  sid 
)    READ_AND_SUM_ST_SPINOR_HALF_TEX_(spinor,sid)

Definition at line 917 of file io_spinor.h.

◆ READ_AND_SUM_ST_SPINOR_HALF_TEX_

#define READ_AND_SUM_ST_SPINOR_HALF_TEX_ (   spinor,
  sid 
)
Value:
{ \
float2 tmp0 = TEX1DFETCH(float2, (spinor), sid + 0*param.sp_stride); \
float2 tmp1 = TEX1DFETCH(float2, (spinor), sid + 1*param.sp_stride); \
float2 tmp2 = TEX1DFETCH(float2, (spinor), sid + 2*param.sp_stride); \
float C = TEX1DFETCH(float, (spinor##Norm), sid); \
o00_re += C*tmp0.x; o00_im += C*tmp0.y; \
o01_re += C*tmp1.x; o01_im += C*tmp1.y; \
o02_re += C*tmp2.x; o02_im += C*tmp2.y; }
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
tmp2
#define tmp2
Definition: tmc_core.h:16
tmp1
#define tmp1
Definition: tmc_core.h:15
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359
tmp0
#define tmp0
Definition: tmc_core.h:14
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 908 of file io_spinor.h.

◆ READ_AND_SUM_ST_SPINOR_SINGLE_TEX

#define READ_AND_SUM_ST_SPINOR_SINGLE_TEX (   spinor,
  sid 
)
Value:
{ \
float2 tmp0 = TEX1DFETCH(float2, (spinor), sid + 0*(param.sp_stride)); \
float2 tmp1 = TEX1DFETCH(float2, (spinor), sid + 1*(param.sp_stride)); \
float2 tmp2 = TEX1DFETCH(float2, (spinor), sid + 2*(param.sp_stride)); \
o00_re += tmp0.x; o00_im += tmp0.y; \
o01_re += tmp1.x; o01_im += tmp1.y; \
o02_re += tmp2.x; o02_im += tmp2.y; }
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
tmp2
#define tmp2
Definition: tmc_core.h:16
tmp1
#define tmp1
Definition: tmc_core.h:15
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359
tmp0
#define tmp0
Definition: tmc_core.h:14
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 900 of file io_spinor.h.

◆ READ_HALF_SPINOR

#define READ_HALF_SPINOR   READ_SPINOR_UP

Definition at line 483 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_DOUBLE

#define READ_KS_NBR_SPINOR_DOUBLE (   T,
  spinor,
  idx,
  mystride 
)
Value:
T##0 = spinor[idx + 0*mystride]; \
T##1 = spinor[idx + 1*mystride]; \
T##2 = spinor[idx + 2*mystride];
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 799 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_DOUBLE_TEX

#define READ_KS_NBR_SPINOR_DOUBLE_TEX (   T,
  spinor,
  idx,
  mystride 
)
Value:
T##0 = fetch_double2((spinor), idx + 0*mystride); \
T##1 = fetch_double2((spinor), idx + 1*mystride); \
T##2 = fetch_double2((spinor), idx + 2*mystride);
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 662 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_GHOST_DOUBLE_TEX

#define READ_KS_NBR_SPINOR_GHOST_DOUBLE_TEX (   T,
  spinor,
  idx,
  mystride,
  dir 
)
Value:
T##0 = fetch_double2((spinor), idx + 0*mystride); \
T##1 = fetch_double2((spinor), idx + 1*mystride); \
T##2 = fetch_double2((spinor), idx + 2*mystride);
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 673 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_GHOST_HALF_TEX

#define READ_KS_NBR_SPINOR_GHOST_HALF_TEX (   T,
  spinor,
  idx,
  mystride,
  dir 
)    READ_KS_NBR_SPINOR_GHOST_HALF_TEX_(T, spinor, idx, mystride, dir)

Definition at line 786 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_GHOST_HALF_TEX_

#define READ_KS_NBR_SPINOR_GHOST_HALF_TEX_ (   T,
  spinor,
  idx,
  mystride,
  dir 
)
Value:
T##0 = TEX1DFETCH(float2, (spinor), idx + 0*mystride); \
T##1 = TEX1DFETCH(float2, (spinor), idx + 1*mystride); \
T##2 = TEX1DFETCH(float2, (spinor), idx + 2*mystride); \
{ \
float C = TEX1DFETCH(float, (spinor ## Norm), norm_idx3); \
(T##0).x *= C; (T##0).y *= C; \
(T##1).x *= C; (T##1).y *= C; \
(T##2).x *= C; (T##2).y *= C;}
x
p x
Definition: CMakeCUDACompilerId.cpp1.ii:3011
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7
y
int y
Definition: CMakeCUDACompilerId.cpp1.ii:2637

Definition at line 772 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_GHOST_SINGLE_TEX

#define READ_KS_NBR_SPINOR_GHOST_SINGLE_TEX (   T,
  spinor,
  idx,
  mystride,
  dir 
)
Value:
T##0 = TEX1DFETCH(float2, (spinor), idx + 0*mystride); \
T##1 = TEX1DFETCH(float2, (spinor), idx + 1*mystride); \
T##2 = TEX1DFETCH(float2, (spinor), idx + 2*mystride);
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 707 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_HALF

#define READ_KS_NBR_SPINOR_HALF (   T,
  spinor,
  idx,
  mystride 
)
Value:
{ \
short2 S0 = in[idx + 0*mystride]; \
short2 S1 = in[idx + 1*mystride]; \
short2 S2 = in[idx + 2*mystride]; \
float C = inNorm[idx]; \
(T##0).x =C*short2float(S0.x); (T##0).y =C*short2float(S0.y); \
(T##1).x =C*short2float(S1.x); (T##1).y =C*short2float(S1.y); \
(T##2).x =C*short2float(S2.x); (T##2).y =C*short2float(S2.y); \
}
x
p x
Definition: CMakeCUDACompilerId.cpp1.ii:3011
in
cpuColorSpinorField * in
Definition: staggered_invert_test.cpp:44
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
y
int y
Definition: CMakeCUDACompilerId.cpp1.ii:2637

Definition at line 831 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_HALF_TEX

#define READ_KS_NBR_SPINOR_HALF_TEX (   T,
  spinor,
  idx,
  mystride 
)    READ_KS_NBR_SPINOR_HALF_TEX_(T, spinor, idx, mystride)

Definition at line 783 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_HALF_TEX_

#define READ_KS_NBR_SPINOR_HALF_TEX_ (   T,
  spinor,
  idx,
  mystride 
)
Value:
T##0 = TEX1DFETCH(float2, (spinor), idx + 0*mystride); \
T##1 = TEX1DFETCH(float2, (spinor), idx + 1*mystride); \
T##2 = TEX1DFETCH(float2, (spinor), idx + 2*mystride); \
{ \
float C = TEX1DFETCH(float, (spinor ## Norm), norm_idx3); \
(T##0).x *= C; (T##0).y *= C; \
(T##1).x *= C; (T##1).y *= C; \
(T##2).x *= C; (T##2).y *= C;}
x
p x
Definition: CMakeCUDACompilerId.cpp1.ii:3011
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7
y
int y
Definition: CMakeCUDACompilerId.cpp1.ii:2637

Definition at line 751 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_SINGLE

#define READ_KS_NBR_SPINOR_SINGLE (   T,
  spinor,
  idx,
  mystride 
)
Value:
T##0 = spinor[idx + 0*mystride]; \
T##1 = spinor[idx + 1*mystride]; \
T##2 = spinor[idx + 2*mystride];
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 814 of file io_spinor.h.

◆ READ_KS_NBR_SPINOR_SINGLE_TEX

#define READ_KS_NBR_SPINOR_SINGLE_TEX (   T,
  spinor,
  idx,
  mystride 
)
Value:
T##0 = TEX1DFETCH(float2, (spinor), idx + 0*mystride); \
T##1 = TEX1DFETCH(float2, (spinor), idx + 1*mystride); \
T##2 = TEX1DFETCH(float2, (spinor), idx + 2*mystride);
idx
int idx
Definition: staggered_fused_exterior_dslash_core.h:355
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 696 of file io_spinor.h.

◆ READ_SPINOR_DOUBLE

#define READ_SPINOR_DOUBLE (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
double2 *spinor = (double2*)spinor_; \
double2 I0 = spinor[sp_idx + 0*(stride)]; \
double2 I1 = spinor[sp_idx + 1*(stride)]; \
double2 I2 = spinor[sp_idx + 2*(stride)]; \
double2 I3 = spinor[sp_idx + 3*(stride)]; \
double2 I4 = spinor[sp_idx + 4*(stride)]; \
double2 I5 = spinor[sp_idx + 5*(stride)]; \
double2 I6 = spinor[sp_idx + 6*(stride)]; \
double2 I7 = spinor[sp_idx + 7*(stride)]; \
double2 I8 = spinor[sp_idx + 8*(stride)]; \
double2 I9 = spinor[sp_idx + 9*(stride)]; \
double2 I10 = spinor[sp_idx + 10*(stride)]; \
double2 I11 = spinor[sp_idx + 11*(stride)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 1 of file io_spinor.h.

◆ READ_SPINOR_DOUBLE_DOWN

#define READ_SPINOR_DOUBLE_DOWN (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
double2 *spinor = (double2*)spinor_; \
double2 I6 = spinor[sp_idx + 6*(stride)]; \
double2 I7 = spinor[sp_idx + 7*(stride)]; \
double2 I8 = spinor[sp_idx + 8*(stride)]; \
double2 I9 = spinor[sp_idx + 9*(stride)]; \
double2 I10 = spinor[sp_idx + 10*(stride)]; \
double2 I11 = spinor[sp_idx + 11*(stride)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 34 of file io_spinor.h.

◆ READ_SPINOR_DOUBLE_DOWN_TEX

#define READ_SPINOR_DOUBLE_DOWN_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
double2 I6 = fetch_double2((spinor), sp_idx + 6*(stride)); \
double2 I7 = fetch_double2((spinor), sp_idx + 7*(stride)); \
double2 I8 = fetch_double2((spinor), sp_idx + 8*(stride)); \
double2 I9 = fetch_double2((spinor), sp_idx + 9*(stride)); \
double2 I10 = fetch_double2((spinor), sp_idx + 10*(stride)); \
double2 I11 = fetch_double2((spinor), sp_idx + 11*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 210 of file io_spinor.h.

◆ READ_SPINOR_DOUBLE_TEX

#define READ_SPINOR_DOUBLE_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
double2 I0 = fetch_double2((spinor), sp_idx + 0*(stride)); \
double2 I1 = fetch_double2((spinor), sp_idx + 1*(stride)); \
double2 I2 = fetch_double2((spinor), sp_idx + 2*(stride)); \
double2 I3 = fetch_double2((spinor), sp_idx + 3*(stride)); \
double2 I4 = fetch_double2((spinor), sp_idx + 4*(stride)); \
double2 I5 = fetch_double2((spinor), sp_idx + 5*(stride)); \
double2 I6 = fetch_double2((spinor), sp_idx + 6*(stride)); \
double2 I7 = fetch_double2((spinor), sp_idx + 7*(stride)); \
double2 I8 = fetch_double2((spinor), sp_idx + 8*(stride)); \
double2 I9 = fetch_double2((spinor), sp_idx + 9*(stride)); \
double2 I10 = fetch_double2((spinor), sp_idx + 10*(stride)); \
double2 I11 = fetch_double2((spinor), sp_idx + 11*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 170 of file io_spinor.h.

◆ READ_SPINOR_DOUBLE_UP

#define READ_SPINOR_DOUBLE_UP (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
double2 *spinor = (double2*)spinor_; \
double2 I0 = spinor[sp_idx + 0*(stride)]; \
double2 I1 = spinor[sp_idx + 1*(stride)]; \
double2 I2 = spinor[sp_idx + 2*(stride)]; \
double2 I3 = spinor[sp_idx + 3*(stride)]; \
double2 I4 = spinor[sp_idx + 4*(stride)]; \
double2 I5 = spinor[sp_idx + 5*(stride)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 25 of file io_spinor.h.

◆ READ_SPINOR_DOUBLE_UP_TEX

#define READ_SPINOR_DOUBLE_UP_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
double2 I0 = fetch_double2((spinor), sp_idx + 0*(stride)); \
double2 I1 = fetch_double2((spinor), sp_idx + 1*(stride)); \
double2 I2 = fetch_double2((spinor), sp_idx + 2*(stride)); \
double2 I3 = fetch_double2((spinor), sp_idx + 3*(stride)); \
double2 I4 = fetch_double2((spinor), sp_idx + 4*(stride)); \
double2 I5 = fetch_double2((spinor), sp_idx + 5*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 202 of file io_spinor.h.

◆ READ_SPINOR_GHOST_DOUBLE

#define READ_SPINOR_GHOST_DOUBLE (   spinor_,
  stride,
  sp_idx,
  norm_idx,
  dir 
)
Value:
double2 *spinor = (double2*)spinor_[dir]; \
double2 I0 = spinor[sp_idx + 0*(stride)]; \
double2 I1 = spinor[sp_idx + 1*(stride)]; \
double2 I2 = spinor[sp_idx + 2*(stride)]; \
double2 I3 = spinor[sp_idx + 3*(stride)]; \
double2 I4 = spinor[sp_idx + 4*(stride)]; \
double2 I5 = spinor[sp_idx + 5*(stride)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 16 of file io_spinor.h.

◆ READ_SPINOR_GHOST_DOUBLE_TEX

#define READ_SPINOR_GHOST_DOUBLE_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx,
  dir 
)
Value:
double2 I0 = fetch_double2((spinor), sp_idx + 0*(stride)); \
double2 I1 = fetch_double2((spinor), sp_idx + 1*(stride)); \
double2 I2 = fetch_double2((spinor), sp_idx + 2*(stride)); \
double2 I3 = fetch_double2((spinor), sp_idx + 3*(stride)); \
double2 I4 = fetch_double2((spinor), sp_idx + 4*(stride)); \
double2 I5 = fetch_double2((spinor), sp_idx + 5*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 193 of file io_spinor.h.

◆ READ_SPINOR_GHOST_HALF

#define READ_SPINOR_GHOST_HALF (   spinor,
  stride,
  sp_idx,
  norm_idx,
  dir 
)    READ_SPINOR_GHOST_HALF_(spinor, stride, sp_idx, norm_idx, dir)

Definition at line 99 of file io_spinor.h.

◆ READ_SPINOR_GHOST_HALF_

#define READ_SPINOR_GHOST_HALF_ (   spinor_,
  stride,
  sp_idx,
  norm_idx,
  dir 
)
Value:
short4 *spinor = (short4*)spinor_[dir]; \
float4 I0 = short42float4(spinor[sp_idx + 0*(stride)]); \
float4 I1 = short42float4(spinor[sp_idx + 1*(stride)]); \
float4 I2 = short42float4(spinor[sp_idx + 2*(stride)]); \
float C = (spinor_ ## Norm)[norm_idx]; \
I0.x *= C; I0.y *= C; I0.z *= C; I0.w *= C; \
I1.x *= C; I1.y *= C; I1.z *= C; I1.w *= C; \
I2.x *= C; I2.y *= C; I2.z *= C; I2.w *= C; \
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 89 of file io_spinor.h.

◆ READ_SPINOR_GHOST_HALF_TEX

#define READ_SPINOR_GHOST_HALF_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx,
  dir 
)    READ_SPINOR_GHOST_HALF_TEX_(spinor, stride, sp_idx, norm_idx, dir) \

Definition at line 300 of file io_spinor.h.

◆ READ_SPINOR_GHOST_HALF_TEX_

#define READ_SPINOR_GHOST_HALF_TEX_ (   spinor,
  stride,
  sp_idx,
  norm_idx,
  dir 
)
Value:
float4 I0 = TEX1DFETCH(float4, (spinor), sp_idx + 0*(stride)); \
float4 I1 = TEX1DFETCH(float4, (spinor), sp_idx + 1*(stride)); \
float4 I2 = TEX1DFETCH(float4, (spinor), sp_idx + 2*(stride)); \
float C = TEX1DFETCH(float, (spinor ## Norm), norm_idx); \
I0.x *= C; I0.y *= C; I0.z *= C; I0.w *= C; \
I1.x *= C; I1.y *= C; I1.z *= C; I1.w *= C; \
I2.x *= C; I2.y *= C; I2.z *= C; I2.w *= C;
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 290 of file io_spinor.h.

◆ READ_SPINOR_GHOST_SINGLE

#define READ_SPINOR_GHOST_SINGLE (   spinor_,
  stride,
  sp_idx,
  norm_idx,
  dir 
)
Value:
float4 *spinor = (float4*)spinor_[dir]; \
float4 I0 = spinor[sp_idx + 0*(stride)]; \
float4 I1 = spinor[sp_idx + 1*(stride)]; \
float4 I2 = spinor[sp_idx + 2*(stride)]; \
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 52 of file io_spinor.h.

◆ READ_SPINOR_GHOST_SINGLE_TEX

#define READ_SPINOR_GHOST_SINGLE_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx,
  dir 
)
Value:
float4 I0 = TEX1DFETCH(float4, (spinor), sp_idx + 0*(stride)); \
float4 I1 = TEX1DFETCH(float4, (spinor), sp_idx + 1*(stride)); \
float4 I2 = TEX1DFETCH(float4, (spinor), sp_idx + 2*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 246 of file io_spinor.h.

◆ READ_SPINOR_HALF

#define READ_SPINOR_HALF (   spinor,
  stride,
  sp_idx,
  norm_idx 
)    READ_SPINOR_HALF_(spinor, stride, sp_idx, norm_idx)

Definition at line 86 of file io_spinor.h.

◆ READ_SPINOR_HALF_

#define READ_SPINOR_HALF_ (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
short4 *spinor = (short4*)spinor_; \
float4 I0 = short42float4(spinor[sp_idx + 0*(stride)]); \
float4 I1 = short42float4(spinor[sp_idx + 1*(stride)]); \
float4 I2 = short42float4(spinor[sp_idx + 2*(stride)]); \
float4 I3 = short42float4(spinor[sp_idx + 3*(stride)]); \
float4 I4 = short42float4(spinor[sp_idx + 4*(stride)]); \
float4 I5 = short42float4(spinor[sp_idx + 5*(stride)]); \
float C = (spinor_ ## Norm)[norm_idx]; \
I0.x *= C; I0.y *= C; I0.z *= C; I0.w *= C; \
I1.x *= C; I1.y *= C; I1.z *= C; I1.w *= C; \
I2.x *= C; I2.y *= C; I2.z *= C; I2.w *= C; \
I3.x *= C; I3.y *= C; I3.z *= C; I3.w *= C; \
I4.x *= C; I4.y *= C; I4.z *= C; I4.w *= C; \
I5.x *= C; I5.y *= C; I5.z *= C; I5.w *= C;
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 70 of file io_spinor.h.

◆ READ_SPINOR_HALF_DOWN

#define READ_SPINOR_HALF_DOWN (   spinor,
  stride,
  sp_idx,
  norm_idx 
)    READ_SPINOR_HALF_DOWN_(spinor, stride, sp_idx, norm_idx)

Definition at line 125 of file io_spinor.h.

◆ READ_SPINOR_HALF_DOWN_

#define READ_SPINOR_HALF_DOWN_ (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
short4 *spinor = (short4*)spinor_; \
float4 I3 = short42float4(spinor[sp_idx + 3*stride]); \
float4 I4 = short42float4(spinor[sp_idx + 4*stride]); \
float4 I5 = short42float4(spinor[sp_idx + 5*stride]); \
float C = (spinor_ ## Norm)[norm_idx]; \
I3.x *= C; I3.y *= C; I3.z *= C; I3.w *= C; \
I4.x *= C; I4.y *= C; I4.z *= C; I4.w *= C; \
I5.x *= C; I5.y *= C; I5.z *= C; I5.w *= C;
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 115 of file io_spinor.h.

◆ READ_SPINOR_HALF_DOWN_TEX

#define READ_SPINOR_HALF_DOWN_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)    READ_SPINOR_HALF_DOWN_TEX_(spinor, stride, sp_idx, norm_idx) \

Definition at line 324 of file io_spinor.h.

◆ READ_SPINOR_HALF_DOWN_TEX_

#define READ_SPINOR_HALF_DOWN_TEX_ (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I3 = TEX1DFETCH(float4, (spinor), sp_idx + 3*(stride)); \
float4 I4 = TEX1DFETCH(float4, (spinor), sp_idx + 4*(stride)); \
float4 I5 = TEX1DFETCH(float4, (spinor), sp_idx + 5*(stride)); \
float C = TEX1DFETCH(float, (spinor ## Norm), norm_idx); \
I3.x *= C; I3.y *= C; I3.z *= C; I3.w *= C; \
I4.x *= C; I4.y *= C; I4.z *= C; I4.w *= C; \
I5.x *= C; I5.y *= C; I5.z *= C; I5.w *= C;
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 315 of file io_spinor.h.

◆ READ_SPINOR_HALF_TEX

#define READ_SPINOR_HALF_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)    READ_SPINOR_HALF_TEX_(spinor, stride, sp_idx, norm_idx) \

Definition at line 277 of file io_spinor.h.

◆ READ_SPINOR_HALF_TEX_

#define READ_SPINOR_HALF_TEX_ (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I0 = TEX1DFETCH(float4, (spinor), sp_idx + 0*(stride)); \
float4 I1 = TEX1DFETCH(float4, (spinor), sp_idx + 1*(stride)); \
float4 I2 = TEX1DFETCH(float4, (spinor), sp_idx + 2*(stride)); \
float4 I3 = TEX1DFETCH(float4, (spinor), sp_idx + 3*(stride)); \
float4 I4 = TEX1DFETCH(float4, (spinor), sp_idx + 4*(stride)); \
float4 I5 = TEX1DFETCH(float4, (spinor), sp_idx + 5*(stride)); \
float C = TEX1DFETCH(float, (spinor ## Norm), norm_idx); \
I0.x *= C; I0.y *= C; I0.z *= C; I0.w *= C; \
I1.x *= C; I1.y *= C; I1.z *= C; I1.w *= C; \
I2.x *= C; I2.y *= C; I2.z *= C; I2.w *= C; \
I3.x *= C; I3.y *= C; I3.z *= C; I3.w *= C; \
I4.x *= C; I4.y *= C; I4.z *= C; I4.w *= C; \
I5.x *= C; I5.y *= C; I5.z *= C; I5.w *= C;
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 262 of file io_spinor.h.

◆ READ_SPINOR_HALF_UP

#define READ_SPINOR_HALF_UP (   spinor,
  stride,
  sp_idx,
  norm_idx 
)    READ_SPINOR_HALF_UP_(spinor, stride, sp_idx, norm_idx)

Definition at line 112 of file io_spinor.h.

◆ READ_SPINOR_HALF_UP_

#define READ_SPINOR_HALF_UP_ (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
short4 *spinor = (short4*)spinor_; \
float4 I0 = short42float4(spinor[sp_idx + 0*(stride)]); \
float4 I1 = short42float4(spinor[sp_idx + 1*(stride)]); \
float4 I2 = short42float4(spinor[sp_idx + 2*(stride)]); \
float C = (spinor_ ## Norm)[norm_idx]; \
I0.x *= C; I0.y *= C; I0.z *= C; I0.w *= C; \
I1.x *= C; I1.y *= C; I1.z *= C; I1.w *= C; \
I2.x *= C; I2.y *= C; I2.z *= C; I2.w *= C; \
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 102 of file io_spinor.h.

◆ READ_SPINOR_HALF_UP_TEX

#define READ_SPINOR_HALF_UP_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)    READ_SPINOR_HALF_UP_TEX_(spinor, stride, sp_idx, norm_idx) \

Definition at line 312 of file io_spinor.h.

◆ READ_SPINOR_HALF_UP_TEX_

#define READ_SPINOR_HALF_UP_TEX_ (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I0 = TEX1DFETCH(float4, (spinor), sp_idx + 0*(stride)); \
float4 I1 = TEX1DFETCH(float4, (spinor), sp_idx + 1*(stride)); \
float4 I2 = TEX1DFETCH(float4, (spinor), sp_idx + 2*(stride)); \
float C = TEX1DFETCH(float, (spinor ## Norm), norm_idx); \
I0.x *= C; I0.y *= C; I0.z *= C; I0.w *= C; \
I1.x *= C; I1.y *= C; I1.z *= C; I1.w *= C; \
I2.x *= C; I2.y *= C; I2.z *= C; I2.w *= C; \
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 303 of file io_spinor.h.

◆ READ_SPINOR_SHARED_DOUBLE2

#define READ_SPINOR_SHARED_DOUBLE2 (   tx,
  ty,
  tz 
)
Value:
extern __shared__ char s_data[]; \
double *sh = (double*)s_data + DSLASH_SHARED_FLOATS_PER_THREAD*SHARED_STRIDE* \
((tx+blockDim.x*(ty+blockDim.y*tz)) / SHARED_STRIDE) + ((tx+blockDim.x*(ty+blockDim.y*tz)) % SHARED_STRIDE); \
double2 I0 = make_double2(sh[0*SHARED_STRIDE], sh[1*SHARED_STRIDE]); \
double2 I1 = make_double2(sh[2*SHARED_STRIDE], sh[3*SHARED_STRIDE]); \
double2 I2 = make_double2(sh[4*SHARED_STRIDE], sh[5*SHARED_STRIDE]); \
double2 I3 = make_double2(sh[6*SHARED_STRIDE], sh[7*SHARED_STRIDE]); \
double2 I4 = make_double2(sh[8*SHARED_STRIDE], sh[9*SHARED_STRIDE]); \
double2 I5 = make_double2(sh[10*SHARED_STRIDE], sh[11*SHARED_STRIDE]); \
double2 I6 = make_double2(sh[12*SHARED_STRIDE], sh[13*SHARED_STRIDE]); \
double2 I7 = make_double2(sh[14*SHARED_STRIDE], sh[15*SHARED_STRIDE]); \
double2 I8 = make_double2(sh[16*SHARED_STRIDE], sh[17*SHARED_STRIDE]); \
double2 I9 = make_double2(sh[18*SHARED_STRIDE], sh[19*SHARED_STRIDE]); \
double2 I10 = make_double2(sh[20*SHARED_STRIDE], sh[21*SHARED_STRIDE]); \
double2 I11 = make_double2(sh[22*SHARED_STRIDE], sh[23*SHARED_STRIDE]);
blockDim
dim3 dim3 blockDim
Definition: CMakeCUDACompilerId.cpp1.ii:2471
tz
int tz
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1810
DSLASH_SHARED_FLOATS_PER_THREAD
#define DSLASH_SHARED_FLOATS_PER_THREAD
Definition: dslash_clover.cu:51
SHARED_STRIDE
#define SHARED_STRIDE
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:386
ty
int ty
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1366
tx
int tx
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1146

Definition at line 1011 of file io_spinor.h.

◆ READ_SPINOR_SHARED_FLOAT4

#define READ_SPINOR_SHARED_FLOAT4 (   tx,
  ty,
  tz 
)
Value:
extern __shared__ char s_data[]; \
float *sh = (float*)s_data + DSLASH_SHARED_FLOATS_PER_THREAD*SHARED_STRIDE* \
((tx+blockDim.x*(ty+blockDim.y*tz)) / SHARED_STRIDE) + ((tx+blockDim.x*(ty+blockDim.y*tz)) % SHARED_STRIDE); \
float4 I0 = make_float4(sh[0*SHARED_STRIDE], sh[1*SHARED_STRIDE], sh[2*SHARED_STRIDE], sh[3*SHARED_STRIDE]); \
float4 I1 = make_float4(sh[4*SHARED_STRIDE], sh[5*SHARED_STRIDE], sh[6*SHARED_STRIDE], sh[7*SHARED_STRIDE]); \
float4 I2 = make_float4(sh[8*SHARED_STRIDE], sh[9*SHARED_STRIDE], sh[10*SHARED_STRIDE], sh[11*SHARED_STRIDE]); \
float4 I3 = make_float4(sh[12*SHARED_STRIDE], sh[13*SHARED_STRIDE], sh[14*SHARED_STRIDE], sh[15*SHARED_STRIDE]); \
float4 I4 = make_float4(sh[16*SHARED_STRIDE], sh[17*SHARED_STRIDE], sh[18*SHARED_STRIDE], sh[19*SHARED_STRIDE]); \
float4 I5 = make_float4(sh[20*SHARED_STRIDE], sh[21*SHARED_STRIDE], sh[22*SHARED_STRIDE], sh[23*SHARED_STRIDE]);
blockDim
dim3 dim3 blockDim
Definition: CMakeCUDACompilerId.cpp1.ii:2471
tz
int tz
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1810
DSLASH_SHARED_FLOATS_PER_THREAD
#define DSLASH_SHARED_FLOATS_PER_THREAD
Definition: dslash_clover.cu:51
SHARED_STRIDE
#define SHARED_STRIDE
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:386
ty
int ty
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1366
tx
int tx
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1146

Definition at line 1032 of file io_spinor.h.

◆ READ_SPINOR_SINGLE

#define READ_SPINOR_SINGLE (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 *spinor = (float4*)spinor_; \
float4 I0 = spinor[sp_idx + 0*(stride)]; \
float4 I1 = spinor[sp_idx + 1*(stride)]; \
float4 I2 = spinor[sp_idx + 2*(stride)]; \
float4 I3 = spinor[sp_idx + 3*(stride)]; \
float4 I4 = spinor[sp_idx + 4*(stride)]; \
float4 I5 = spinor[sp_idx + 5*(stride)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 43 of file io_spinor.h.

◆ READ_SPINOR_SINGLE_DOWN

#define READ_SPINOR_SINGLE_DOWN (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 *spinor = (float4*)spinor_; \
float4 I3 = spinor[sp_idx + 3*(stride)]; \
float4 I4 = spinor[sp_idx + 4*(stride)]; \
float4 I5 = spinor[sp_idx + 5*(stride)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 64 of file io_spinor.h.

◆ READ_SPINOR_SINGLE_DOWN_TEX

#define READ_SPINOR_SINGLE_DOWN_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I3 = TEX1DFETCH(float4, (spinor), sp_idx + 3*(stride)); \
float4 I4 = TEX1DFETCH(float4, (spinor), sp_idx + 4*(stride)); \
float4 I5 = TEX1DFETCH(float4, (spinor), sp_idx + 5*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 257 of file io_spinor.h.

◆ READ_SPINOR_SINGLE_TEX

#define READ_SPINOR_SINGLE_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I0 = TEX1DFETCH(float4, (spinor), sp_idx + 0*(stride)); \
float4 I1 = TEX1DFETCH(float4, (spinor), sp_idx + 1*(stride)); \
float4 I2 = TEX1DFETCH(float4, (spinor), sp_idx + 2*(stride)); \
float4 I3 = TEX1DFETCH(float4, (spinor), sp_idx + 3*(stride)); \
float4 I4 = TEX1DFETCH(float4, (spinor), sp_idx + 4*(stride)); \
float4 I5 = TEX1DFETCH(float4, (spinor), sp_idx + 5*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 232 of file io_spinor.h.

◆ READ_SPINOR_SINGLE_UP

#define READ_SPINOR_SINGLE_UP (   spinor_,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 *spinor = (float4*)spinor_; \
float4 I0 = spinor[sp_idx + 0*(stride)]; \
float4 I1 = spinor[sp_idx + 1*(stride)]; \
float4 I2 = spinor[sp_idx + 2*(stride)]; \
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 58 of file io_spinor.h.

◆ READ_SPINOR_SINGLE_UP_TEX

#define READ_SPINOR_SINGLE_UP_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I0 = TEX1DFETCH(float4, (spinor), sp_idx + 0*(stride)); \
float4 I1 = TEX1DFETCH(float4, (spinor), sp_idx + 1*(stride)); \
float4 I2 = TEX1DFETCH(float4, (spinor), sp_idx + 2*(stride)); \
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 252 of file io_spinor.h.

◆ READ_ST_ACCUM_DOUBLE

#define READ_ST_ACCUM_DOUBLE (   spinor,
  sid 
)
Value:
double2 accum0 = spinor[sid + 0*(param.sp_stride)]; \
double2 accum1 = spinor[sid + 1*(param.sp_stride)]; \
double2 accum2 = spinor[sid + 2*(param.sp_stride)];
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
param
QudaGaugeParam param
Definition: pack_test.cpp:17
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 958 of file io_spinor.h.

◆ READ_ST_ACCUM_DOUBLE_TEX

#define READ_ST_ACCUM_DOUBLE_TEX (   spinor,
  sid 
)
Value:
double2 accum0 = fetch_double2((spinor), sid + 0*(param.sp_stride)); \
double2 accum1 = fetch_double2((spinor), sid + 1*(param.sp_stride)); \
double2 accum2 = fetch_double2((spinor), sid + 2*(param.sp_stride));
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
param
QudaGaugeParam param
Definition: pack_test.cpp:17
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 937 of file io_spinor.h.

◆ READ_ST_ACCUM_HALF

#define READ_ST_ACCUM_HALF (   spinor,
  sid 
)
Value:
float2 accum0, accum1, accum2; \
{ \
short2 S0 = x[sid + 0*param.sp_stride]; \
short2 S1 = x[sid + 1*param.sp_stride]; \
short2 S2 = x[sid + 2*param.sp_stride]; \
float C = spinor##Norm[sid]; \
accum0.x =C*short2float(S0.x); accum0.y =C*short2float(S0.y); \
accum1.x =C*short2float(S1.x); accum1.y =C*short2float(S1.y); \
accum2.x =C*short2float(S2.x); accum2.y =C*short2float(S2.y); \
}
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
param
QudaGaugeParam param
Definition: pack_test.cpp:17
x
p x
Definition: CMakeCUDACompilerId.cpp1.ii:3011
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 968 of file io_spinor.h.

◆ READ_ST_ACCUM_HALF_TEX

#define READ_ST_ACCUM_HALF_TEX (   spinor,
  sid 
)    READ_ST_ACCUM_HALF_TEX_(spinor,sid)

Definition at line 956 of file io_spinor.h.

◆ READ_ST_ACCUM_HALF_TEX_

#define READ_ST_ACCUM_HALF_TEX_ (   spinor,
  sid 
)
Value:
float2 accum0 = TEX1DFETCH(float2, (spinor), sid + 0*param.sp_stride); \
float2 accum1 = TEX1DFETCH(float2, (spinor), sid + 1*param.sp_stride); \
float2 accum2 = TEX1DFETCH(float2, (spinor), sid + 2*param.sp_stride); \
float C = TEX1DFETCH(float, (spinor ## Norm), sid); \
accum0.x *= C; accum0.y *= C; \
accum1.x *= C; accum1.y *= C; \
accum2.x *= C; accum2.y *= C;
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
param
QudaGaugeParam param
Definition: pack_test.cpp:17
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 947 of file io_spinor.h.

◆ READ_ST_ACCUM_SINGLE

#define READ_ST_ACCUM_SINGLE (   spinor,
  sid 
)
Value:
float2 accum0 = spinor[sid + 0*(param.sp_stride)]; \
float2 accum1 = spinor[sid + 1*(param.sp_stride)]; \
float2 accum2 = spinor[sid + 2*(param.sp_stride)];
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
param
QudaGaugeParam param
Definition: pack_test.cpp:17
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 963 of file io_spinor.h.

◆ READ_ST_ACCUM_SINGLE_TEX

#define READ_ST_ACCUM_SINGLE_TEX (   spinor,
  sid 
)
Value:
float2 accum0 = TEX1DFETCH(float2, (spinor), sid + 0*param.sp_stride); \
float2 accum1 = TEX1DFETCH(float2, (spinor), sid + 1*param.sp_stride); \
float2 accum2 = TEX1DFETCH(float2, (spinor), sid + 2*param.sp_stride);
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
param
QudaGaugeParam param
Definition: pack_test.cpp:17
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 942 of file io_spinor.h.

◆ WRITE_FLAVOR_SPINOR_DOUBLE2

#define WRITE_FLAVOR_SPINOR_DOUBLE2 ( )
Value:
double2 *out = (double2*)param.out; \
out[0*(param.sp_stride)+sid] = make_double2(o1_00_re, o1_00_im); \
out[1*(param.sp_stride)+sid] = make_double2(o1_01_re, o1_01_im); \
out[2*(param.sp_stride)+sid] = make_double2(o1_02_re, o1_02_im); \
out[3*(param.sp_stride)+sid] = make_double2(o1_10_re, o1_10_im); \
out[4*(param.sp_stride)+sid] = make_double2(o1_11_re, o1_11_im); \
out[5*(param.sp_stride)+sid] = make_double2(o1_12_re, o1_12_im); \
out[6*(param.sp_stride)+sid] = make_double2(o1_20_re, o1_20_im); \
out[7*(param.sp_stride)+sid] = make_double2(o1_21_re, o1_21_im); \
out[8*(param.sp_stride)+sid] = make_double2(o1_22_re, o1_22_im); \
out[9*(param.sp_stride)+sid] = make_double2(o1_30_re, o1_30_im); \
out[10*(param.sp_stride)+sid] = make_double2(o1_31_re, o1_31_im); \
out[11*(param.sp_stride)+sid] = make_double2(o1_32_re, o1_32_im); \
out[0*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_00_re, o2_00_im); \
out[1*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_01_re, o2_01_im); \
out[2*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_02_re, o2_02_im); \
out[3*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_10_re, o2_10_im); \
out[4*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_11_re, o2_11_im); \
out[5*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_12_re, o2_12_im); \
out[6*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_20_re, o2_20_im); \
out[7*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_21_re, o2_21_im); \
out[8*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_22_re, o2_22_im); \
out[9*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_30_re, o2_30_im); \
out[10*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_31_re, o2_31_im); \
out[11*(param.sp_stride)+sid+param.fl_stride] = make_double2(o2_32_re, o2_32_im);
o2_01_re
VOLATILE spinorFloat o2_01_re
Definition: tm_ndeg_dslash_core.h:159
o1_30_re
VOLATILE spinorFloat o1_30_re
Definition: tm_ndeg_dslash_core.h:150
o1_02_im
VOLATILE spinorFloat o1_02_im
Definition: tm_ndeg_dslash_core.h:137
o1_32_re
VOLATILE spinorFloat o1_32_re
Definition: tm_ndeg_dslash_core.h:154
o1_22_im
VOLATILE spinorFloat o1_22_im
Definition: tm_ndeg_dslash_core.h:149
o2_32_im
VOLATILE spinorFloat o2_32_im
Definition: tm_ndeg_dslash_core.h:180
o2_02_im
VOLATILE spinorFloat o2_02_im
Definition: tm_ndeg_dslash_core.h:162
o2_31_re
VOLATILE spinorFloat o2_31_re
Definition: tm_ndeg_dslash_core.h:177
o1_31_re
VOLATILE spinorFloat o1_31_re
Definition: tm_ndeg_dslash_core.h:152
o2_32_re
VOLATILE spinorFloat o2_32_re
Definition: tm_ndeg_dslash_core.h:179
o1_20_re
VOLATILE spinorFloat o1_20_re
Definition: tm_ndeg_dslash_core.h:144
o1_02_re
VOLATILE spinorFloat o1_02_re
Definition: tm_ndeg_dslash_core.h:136
o2_22_re
VOLATILE spinorFloat o2_22_re
Definition: tm_ndeg_dslash_core.h:173
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o1_10_im
VOLATILE spinorFloat o1_10_im
Definition: tm_ndeg_dslash_core.h:139
o2_30_re
VOLATILE spinorFloat o2_30_re
Definition: tm_ndeg_dslash_core.h:175
o2_31_im
VOLATILE spinorFloat o2_31_im
Definition: tm_ndeg_dslash_core.h:178
o1_31_im
VOLATILE spinorFloat o1_31_im
Definition: tm_ndeg_dslash_core.h:153
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o1_22_re
VOLATILE spinorFloat o1_22_re
Definition: tm_ndeg_dslash_core.h:148
o1_00_re
VOLATILE spinorFloat o1_00_re
Definition: tm_ndeg_dslash_core.h:132
o1_12_re
VOLATILE spinorFloat o1_12_re
Definition: tm_ndeg_dslash_core.h:142
o2_12_re
VOLATILE spinorFloat o2_12_re
Definition: tm_ndeg_dslash_core.h:167
o1_11_re
VOLATILE spinorFloat o1_11_re
Definition: tm_ndeg_dslash_core.h:140
o2_22_im
VOLATILE spinorFloat o2_22_im
Definition: tm_ndeg_dslash_core.h:174
o2_00_im
VOLATILE spinorFloat o2_00_im
Definition: tm_ndeg_dslash_core.h:158
o2_21_im
VOLATILE spinorFloat o2_21_im
Definition: tm_ndeg_dslash_core.h:172
o1_32_im
VOLATILE spinorFloat o1_32_im
Definition: tm_ndeg_dslash_core.h:155
o1_21_im
VOLATILE spinorFloat o1_21_im
Definition: tm_ndeg_dslash_core.h:147
o2_10_im
VOLATILE spinorFloat o2_10_im
Definition: tm_ndeg_dslash_core.h:164
o1_01_re
VOLATILE spinorFloat o1_01_re
Definition: tm_ndeg_dslash_core.h:134
o2_01_im
VOLATILE spinorFloat o2_01_im
Definition: tm_ndeg_dslash_core.h:160
o1_01_im
VOLATILE spinorFloat o1_01_im
Definition: tm_ndeg_dslash_core.h:135
o1_11_im
VOLATILE spinorFloat o1_11_im
Definition: tm_ndeg_dslash_core.h:141
o1_12_im
VOLATILE spinorFloat o1_12_im
Definition: tm_ndeg_dslash_core.h:143
o2_10_re
VOLATILE spinorFloat o2_10_re
Definition: tm_ndeg_dslash_core.h:163
o2_02_re
VOLATILE spinorFloat o2_02_re
Definition: tm_ndeg_dslash_core.h:161
o1_20_im
VOLATILE spinorFloat o1_20_im
Definition: tm_ndeg_dslash_core.h:145
o2_00_re
VOLATILE spinorFloat o2_00_re
Definition: tm_ndeg_dslash_core.h:157
o2_11_im
VOLATILE spinorFloat o2_11_im
Definition: tm_ndeg_dslash_core.h:166
o1_10_re
VOLATILE spinorFloat o1_10_re
Definition: tm_ndeg_dslash_core.h:138
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o1_30_im
VOLATILE spinorFloat o1_30_im
Definition: tm_ndeg_dslash_core.h:151
o2_20_im
VOLATILE spinorFloat o2_20_im
Definition: tm_ndeg_dslash_core.h:170
o2_30_im
VOLATILE spinorFloat o2_30_im
Definition: tm_ndeg_dslash_core.h:176
o2_12_im
VOLATILE spinorFloat o2_12_im
Definition: tm_ndeg_dslash_core.h:168
o2_21_re
VOLATILE spinorFloat o2_21_re
Definition: tm_ndeg_dslash_core.h:171
o1_21_re
VOLATILE spinorFloat o1_21_re
Definition: tm_ndeg_dslash_core.h:146
o2_20_re
VOLATILE spinorFloat o2_20_re
Definition: tm_ndeg_dslash_core.h:169
o2_11_re
VOLATILE spinorFloat o2_11_re
Definition: tm_ndeg_dslash_core.h:165
o1_00_im
VOLATILE spinorFloat o1_00_im
Definition: tm_ndeg_dslash_core.h:133

ndeg tm:

Definition at line 521 of file io_spinor.h.

◆ WRITE_FLAVOR_SPINOR_FLOAT4

#define WRITE_FLAVOR_SPINOR_FLOAT4 ( )
Value:
float4 *out = (float4*)param.out; \
out[0*(param.sp_stride)+sid] = make_float4(o1_00_re, o1_00_im, o1_01_re, o1_01_im); \
out[1*(param.sp_stride)+sid] = make_float4(o1_02_re, o1_02_im, o1_10_re, o1_10_im); \
out[2*(param.sp_stride)+sid] = make_float4(o1_11_re, o1_11_im, o1_12_re, o1_12_im); \
out[3*(param.sp_stride)+sid] = make_float4(o1_20_re, o1_20_im, o1_21_re, o1_21_im); \
out[4*(param.sp_stride)+sid] = make_float4(o1_22_re, o1_22_im, o1_30_re, o1_30_im); \
out[5*(param.sp_stride)+sid] = make_float4(o1_31_re, o1_31_im, o1_32_re, o1_32_im); \
out[0*(param.sp_stride)+sid+param.fl_stride] = make_float4(o2_00_re, o2_00_im, o2_01_re, o2_01_im); \
out[1*(param.sp_stride)+sid+param.fl_stride] = make_float4(o2_02_re, o2_02_im, o2_10_re, o2_10_im); \
out[2*(param.sp_stride)+sid+param.fl_stride] = make_float4(o2_11_re, o2_11_im, o2_12_re, o2_12_im); \
out[3*(param.sp_stride)+sid+param.fl_stride] = make_float4(o2_20_re, o2_20_im, o2_21_re, o2_21_im); \
out[4*(param.sp_stride)+sid+param.fl_stride] = make_float4(o2_22_re, o2_22_im, o2_30_re, o2_30_im); \
out[5*(param.sp_stride)+sid+param.fl_stride] = make_float4(o2_31_re, o2_31_im, o2_32_re, o2_32_im);
o2_01_re
VOLATILE spinorFloat o2_01_re
Definition: tm_ndeg_dslash_core.h:159
o1_30_re
VOLATILE spinorFloat o1_30_re
Definition: tm_ndeg_dslash_core.h:150
o1_02_im
VOLATILE spinorFloat o1_02_im
Definition: tm_ndeg_dslash_core.h:137
o1_32_re
VOLATILE spinorFloat o1_32_re
Definition: tm_ndeg_dslash_core.h:154
o1_22_im
VOLATILE spinorFloat o1_22_im
Definition: tm_ndeg_dslash_core.h:149
o2_32_im
VOLATILE spinorFloat o2_32_im
Definition: tm_ndeg_dslash_core.h:180
o2_02_im
VOLATILE spinorFloat o2_02_im
Definition: tm_ndeg_dslash_core.h:162
o2_31_re
VOLATILE spinorFloat o2_31_re
Definition: tm_ndeg_dslash_core.h:177
o1_31_re
VOLATILE spinorFloat o1_31_re
Definition: tm_ndeg_dslash_core.h:152
o2_32_re
VOLATILE spinorFloat o2_32_re
Definition: tm_ndeg_dslash_core.h:179
o1_20_re
VOLATILE spinorFloat o1_20_re
Definition: tm_ndeg_dslash_core.h:144
o1_02_re
VOLATILE spinorFloat o1_02_re
Definition: tm_ndeg_dslash_core.h:136
o2_22_re
VOLATILE spinorFloat o2_22_re
Definition: tm_ndeg_dslash_core.h:173
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o1_10_im
VOLATILE spinorFloat o1_10_im
Definition: tm_ndeg_dslash_core.h:139
o2_30_re
VOLATILE spinorFloat o2_30_re
Definition: tm_ndeg_dslash_core.h:175
o2_31_im
VOLATILE spinorFloat o2_31_im
Definition: tm_ndeg_dslash_core.h:178
o1_31_im
VOLATILE spinorFloat o1_31_im
Definition: tm_ndeg_dslash_core.h:153
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o1_22_re
VOLATILE spinorFloat o1_22_re
Definition: tm_ndeg_dslash_core.h:148
o1_00_re
VOLATILE spinorFloat o1_00_re
Definition: tm_ndeg_dslash_core.h:132
o1_12_re
VOLATILE spinorFloat o1_12_re
Definition: tm_ndeg_dslash_core.h:142
o2_12_re
VOLATILE spinorFloat o2_12_re
Definition: tm_ndeg_dslash_core.h:167
o1_11_re
VOLATILE spinorFloat o1_11_re
Definition: tm_ndeg_dslash_core.h:140
o2_22_im
VOLATILE spinorFloat o2_22_im
Definition: tm_ndeg_dslash_core.h:174
o2_00_im
VOLATILE spinorFloat o2_00_im
Definition: tm_ndeg_dslash_core.h:158
o2_21_im
VOLATILE spinorFloat o2_21_im
Definition: tm_ndeg_dslash_core.h:172
o1_32_im
VOLATILE spinorFloat o1_32_im
Definition: tm_ndeg_dslash_core.h:155
o1_21_im
VOLATILE spinorFloat o1_21_im
Definition: tm_ndeg_dslash_core.h:147
o2_10_im
VOLATILE spinorFloat o2_10_im
Definition: tm_ndeg_dslash_core.h:164
o1_01_re
VOLATILE spinorFloat o1_01_re
Definition: tm_ndeg_dslash_core.h:134
o2_01_im
VOLATILE spinorFloat o2_01_im
Definition: tm_ndeg_dslash_core.h:160
o1_01_im
VOLATILE spinorFloat o1_01_im
Definition: tm_ndeg_dslash_core.h:135
o1_11_im
VOLATILE spinorFloat o1_11_im
Definition: tm_ndeg_dslash_core.h:141
o1_12_im
VOLATILE spinorFloat o1_12_im
Definition: tm_ndeg_dslash_core.h:143
o2_10_re
VOLATILE spinorFloat o2_10_re
Definition: tm_ndeg_dslash_core.h:163
o2_02_re
VOLATILE spinorFloat o2_02_re
Definition: tm_ndeg_dslash_core.h:161
o1_20_im
VOLATILE spinorFloat o1_20_im
Definition: tm_ndeg_dslash_core.h:145
o2_00_re
VOLATILE spinorFloat o2_00_re
Definition: tm_ndeg_dslash_core.h:157
o2_11_im
VOLATILE spinorFloat o2_11_im
Definition: tm_ndeg_dslash_core.h:166
o1_10_re
VOLATILE spinorFloat o1_10_re
Definition: tm_ndeg_dslash_core.h:138
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o1_30_im
VOLATILE spinorFloat o1_30_im
Definition: tm_ndeg_dslash_core.h:151
o2_20_im
VOLATILE spinorFloat o2_20_im
Definition: tm_ndeg_dslash_core.h:170
o2_30_im
VOLATILE spinorFloat o2_30_im
Definition: tm_ndeg_dslash_core.h:176
o2_12_im
VOLATILE spinorFloat o2_12_im
Definition: tm_ndeg_dslash_core.h:168
o2_21_re
VOLATILE spinorFloat o2_21_re
Definition: tm_ndeg_dslash_core.h:171
o1_21_re
VOLATILE spinorFloat o1_21_re
Definition: tm_ndeg_dslash_core.h:146
o2_20_re
VOLATILE spinorFloat o2_20_re
Definition: tm_ndeg_dslash_core.h:169
o2_11_re
VOLATILE spinorFloat o2_11_re
Definition: tm_ndeg_dslash_core.h:165
o1_00_im
VOLATILE spinorFloat o1_00_im
Definition: tm_ndeg_dslash_core.h:133

Definition at line 549 of file io_spinor.h.

◆ WRITE_FLAVOR_SPINOR_SHORT4

#define WRITE_FLAVOR_SPINOR_SHORT4 ( )

Definition at line 565 of file io_spinor.h.

◆ WRITE_HALF_SPINOR_DOUBLE2

#define WRITE_HALF_SPINOR_DOUBLE2 (   stride,
  sid 
)
Value:
out[0*(stride)+sid] = make_double2(a0_re, a0_im); \
out[1*(stride)+sid] = make_double2(a1_re, a1_im); \
out[2*(stride)+sid] = make_double2(a2_re, a2_im); \
out[3*(stride)+sid] = make_double2(b0_re, b0_im); \
out[4*(stride)+sid] = make_double2(b1_re, b1_im); \
out[5*(stride)+sid] = make_double2(b2_re, b2_im);
a1_re
spinorFloat a1_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:929
a2_im
spinorFloat a2_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:930
b1_im
spinorFloat b1_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:932
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
b2_im
spinorFloat b2_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:933
a0_im
spinorFloat a0_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:928
b1_re
spinorFloat b1_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:932
a1_im
spinorFloat a1_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:929
b0_im
spinorFloat b0_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:931
b0_re
spinorFloat b0_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:931
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
a2_re
spinorFloat a2_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:930
b2_re
spinorFloat b2_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:933
a0_re
spinorFloat a0_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:928

Definition at line 485 of file io_spinor.h.

◆ WRITE_HALF_SPINOR_FLOAT4

#define WRITE_HALF_SPINOR_FLOAT4 (   stride,
  sid 
)
Value:
out[0*(stride)+sid] = make_float4(a0_re, a0_im, a1_re, a1_im); \
out[1*(stride)+sid] = make_float4(a2_re, a2_im, b0_re, b0_im); \
out[2*(stride)+sid] = make_float4(b1_re, b1_im, b2_re, b2_im);
a1_re
spinorFloat a1_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:929
a2_im
spinorFloat a2_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:930
b1_im
spinorFloat b1_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:932
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
b2_im
spinorFloat b2_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:933
a0_im
spinorFloat a0_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:928
b1_re
spinorFloat b1_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:932
a1_im
spinorFloat a1_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:929
b0_im
spinorFloat b0_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:931
b0_re
spinorFloat b0_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:931
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
a2_re
spinorFloat a2_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:930
b2_re
spinorFloat b2_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:933
a0_re
spinorFloat a0_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:928

Definition at line 493 of file io_spinor.h.

◆ WRITE_HALF_SPINOR_SHORT4

#define WRITE_HALF_SPINOR_SHORT4 (   stride,
  sid 
)
Value:
float c0 = fmaxf(fabsf(a0_re), fabsf(a0_im)); \
float c1 = fmaxf(fabsf(a1_re), fabsf(a1_im)); \
float c2 = fmaxf(fabsf(a2_re), fabsf(a2_im)); \
float c3 = fmaxf(fabsf(b0_re), fabsf(b0_im)); \
float c4 = fmaxf(fabsf(b1_re), fabsf(b1_im)); \
float c5 = fmaxf(fabsf(b2_re), fabsf(b2_im)); \
c0 = fmaxf(c0, c1); \
c1 = fmaxf(c2, c3); \
c2 = fmaxf(c4, c5); \
c0 = fmaxf(c0, c1); \
c0 = fmaxf(c0, c2); \
outNorm[sid] = c0; \
float scale = __fdividef(MAX_SHORT, c0); \
a0_re *= scale; a0_im *= scale; a1_re *= scale; a1_im *= scale; \
a2_re *= scale; a2_im *= scale; b0_re *= scale; b0_im *= scale; \
b1_re *= scale; b1_im *= scale; b2_re *= scale; b2_im *= scale; \
out[sid+0*(stride)] = make_short4(f2i(a0_re), f2i(a0_im), f2i(a1_re), f2i(a1_im)); \
out[sid+1*(stride)] = make_short4(f2i(a2_re), f2i(a2_im), f2i(b0_re), f2i(b0_im)); \
out[sid+2*(stride)] = make_short4(f2i(b1_re), f2i(b1_im), f2i(b2_re), f2i(b2_im));
a1_re
spinorFloat a1_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:929
f2i
__device__ int f2i(float f)
Definition: convert.h:47
fabsf
float fabsf(float)
a2_im
spinorFloat a2_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:930
b1_im
spinorFloat b1_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:932
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
b2_im
spinorFloat b2_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:933
a0_im
spinorFloat a0_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:928
b1_re
spinorFloat b1_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:932
fmaxf
float fmaxf(float, float)
a1_im
spinorFloat a1_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:929
b0_im
spinorFloat b0_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:931
b0_re
spinorFloat b0_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:931
a2_re
spinorFloat a2_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:930
MAX_SHORT
#define MAX_SHORT
Definition: quda_internal.h:29
b2_re
spinorFloat b2_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:933
a0_re
spinorFloat a0_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:928

Definition at line 498 of file io_spinor.h.

◆ WRITE_SPINOR_DOUBLE2

#define WRITE_SPINOR_DOUBLE2 (   stride)
Value:
double2 *out = (double2*)param.out; \
out[0*(stride)+sid] = make_double2(o00_re, o00_im); \
out[1*(stride)+sid] = make_double2(o01_re, o01_im); \
out[2*(stride)+sid] = make_double2(o02_re, o02_im); \
out[3*(stride)+sid] = make_double2(o10_re, o10_im); \
out[4*(stride)+sid] = make_double2(o11_re, o11_im); \
out[5*(stride)+sid] = make_double2(o12_re, o12_im); \
out[6*(stride)+sid] = make_double2(o20_re, o20_im); \
out[7*(stride)+sid] = make_double2(o21_re, o21_im); \
out[8*(stride)+sid] = make_double2(o22_re, o22_im); \
out[9*(stride)+sid] = make_double2(o30_re, o30_im); \
out[10*(stride)+sid] = make_double2(o31_re, o31_im); \
out[11*(stride)+sid] = make_double2(o32_re, o32_im);
o22_im
VOLATILE spinorFloat o22_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:375
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o20_re
VOLATILE spinorFloat o20_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:370
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o21_im
VOLATILE spinorFloat o21_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:373
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
o20_im
VOLATILE spinorFloat o20_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:371
o32_im
VOLATILE spinorFloat o32_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:381
o31_im
VOLATILE spinorFloat o31_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:379
o10_im
VOLATILE spinorFloat o10_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:365
o11_im
VOLATILE spinorFloat o11_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:367
o31_re
VOLATILE spinorFloat o31_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:378
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o10_re
VOLATILE spinorFloat o10_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:364
o11_re
VOLATILE spinorFloat o11_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:366
o21_re
VOLATILE spinorFloat o21_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:372
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359
o12_im
VOLATILE spinorFloat o12_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:369
o12_re
VOLATILE spinorFloat o12_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:368
o30_im
VOLATILE spinorFloat o30_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:377
o32_re
VOLATILE spinorFloat o32_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:380
o22_re
VOLATILE spinorFloat o22_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:374
o30_re
VOLATILE spinorFloat o30_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:376

Definition at line 353 of file io_spinor.h.

◆ WRITE_SPINOR_DOUBLE2_STR

#define WRITE_SPINOR_DOUBLE2_STR (   stride)
Value:
double2 *out = (double2*)param.out; \
store_streaming_double2(&out[0*stride+sid], o00_re, o00_im); \
store_streaming_double2(&out[1*stride+sid], o01_re, o01_im); \
store_streaming_double2(&out[2*stride+sid], o02_re, o02_im); \
store_streaming_double2(&out[3*stride+sid], o10_re, o10_im); \
store_streaming_double2(&out[4*stride+sid], o11_re, o11_im); \
store_streaming_double2(&out[5*stride+sid], o12_re, o12_im); \
store_streaming_double2(&out[6*stride+sid], o20_re, o20_im); \
store_streaming_double2(&out[7*stride+sid], o21_re, o21_im); \
store_streaming_double2(&out[8*stride+sid], o22_re, o22_im); \
store_streaming_double2(&out[9*stride+sid], o30_re, o30_im); \
store_streaming_double2(&out[10*stride+sid], o31_re, o31_im); \
store_streaming_double2(&out[11*stride+sid], o32_re, o32_im);
o22_im
VOLATILE spinorFloat o22_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:375
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o20_re
VOLATILE spinorFloat o20_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:370
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o21_im
VOLATILE spinorFloat o21_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:373
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
o20_im
VOLATILE spinorFloat o20_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:371
o32_im
VOLATILE spinorFloat o32_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:381
o31_im
VOLATILE spinorFloat o31_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:379
quda::store_streaming_double2
__device__ void store_streaming_double2(double2 *addr, double x, double y)
Definition: inline_ptx.h:49
o10_im
VOLATILE spinorFloat o10_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:365
o11_im
VOLATILE spinorFloat o11_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:367
o31_re
VOLATILE spinorFloat o31_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:378
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o10_re
VOLATILE spinorFloat o10_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:364
o11_re
VOLATILE spinorFloat o11_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:366
o21_re
VOLATILE spinorFloat o21_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:372
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359
o12_im
VOLATILE spinorFloat o12_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:369
o12_re
VOLATILE spinorFloat o12_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:368
o30_im
VOLATILE spinorFloat o30_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:377
o32_re
VOLATILE spinorFloat o32_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:380
o22_re
VOLATILE spinorFloat o22_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:374
o30_re
VOLATILE spinorFloat o30_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:376

Definition at line 417 of file io_spinor.h.

◆ WRITE_SPINOR_FLOAT4

#define WRITE_SPINOR_FLOAT4 (   stride)
Value:
float4 *out = (float4*)param.out; \
out[0*(stride)+sid] = make_float4(o00_re, o00_im, o01_re, o01_im); \
out[1*(stride)+sid] = make_float4(o02_re, o02_im, o10_re, o10_im); \
out[2*(stride)+sid] = make_float4(o11_re, o11_im, o12_re, o12_im); \
out[3*(stride)+sid] = make_float4(o20_re, o20_im, o21_re, o21_im); \
out[4*(stride)+sid] = make_float4(o22_re, o22_im, o30_re, o30_im); \
out[5*(stride)+sid] = make_float4(o31_re, o31_im, o32_re, o32_im);
o22_im
VOLATILE spinorFloat o22_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:375
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o20_re
VOLATILE spinorFloat o20_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:370
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o21_im
VOLATILE spinorFloat o21_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:373
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
o20_im
VOLATILE spinorFloat o20_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:371
o32_im
VOLATILE spinorFloat o32_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:381
o31_im
VOLATILE spinorFloat o31_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:379
o10_im
VOLATILE spinorFloat o10_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:365
o11_im
VOLATILE spinorFloat o11_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:367
o31_re
VOLATILE spinorFloat o31_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:378
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o10_re
VOLATILE spinorFloat o10_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:364
o11_re
VOLATILE spinorFloat o11_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:366
o21_re
VOLATILE spinorFloat o21_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:372
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359
o12_im
VOLATILE spinorFloat o12_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:369
o12_re
VOLATILE spinorFloat o12_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:368
o30_im
VOLATILE spinorFloat o30_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:377
o32_re
VOLATILE spinorFloat o32_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:380
o22_re
VOLATILE spinorFloat o22_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:374
o30_re
VOLATILE spinorFloat o30_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:376

Definition at line 368 of file io_spinor.h.

◆ WRITE_SPINOR_FLOAT4_STR

#define WRITE_SPINOR_FLOAT4_STR (   stride)
Value:
float4 *out = (float4*)param.out; \
store_streaming_float4(&out[0*(stride)+sid], o00_re, o00_im, o01_re, o01_im); \
store_streaming_float4(&out[1*(stride)+sid], o02_re, o02_im, o10_re, o10_im); \
store_streaming_float4(&out[2*(stride)+sid], o11_re, o11_im, o12_re, o12_im); \
store_streaming_float4(&out[3*(stride)+sid], o20_re, o20_im, o21_re, o21_im); \
store_streaming_float4(&out[4*(stride)+sid], o22_re, o22_im, o30_re, o30_im); \
store_streaming_float4(&out[5*(stride)+sid], o31_re, o31_im, o32_re, o32_im);
o22_im
VOLATILE spinorFloat o22_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:375
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o20_re
VOLATILE spinorFloat o20_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:370
param
QudaGaugeParam param
Definition: pack_test.cpp:17
o21_im
VOLATILE spinorFloat o21_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:373
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
o20_im
VOLATILE spinorFloat o20_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:371
o32_im
VOLATILE spinorFloat o32_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:381
o31_im
VOLATILE spinorFloat o31_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:379
o10_im
VOLATILE spinorFloat o10_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:365
o11_im
VOLATILE spinorFloat o11_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:367
quda::store_streaming_float4
__device__ void store_streaming_float4(float4 *addr, float x, float y, float z, float w)
Definition: inline_ptx.h:39
o31_re
VOLATILE spinorFloat o31_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:378
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o10_re
VOLATILE spinorFloat o10_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:364
o11_re
VOLATILE spinorFloat o11_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:366
o21_re
VOLATILE spinorFloat o21_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:372
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359
o12_im
VOLATILE spinorFloat o12_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:369
o12_re
VOLATILE spinorFloat o12_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:368
o30_im
VOLATILE spinorFloat o30_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:377
o32_re
VOLATILE spinorFloat o32_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:380
o22_re
VOLATILE spinorFloat o22_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:374
o30_re
VOLATILE spinorFloat o30_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:376

Definition at line 432 of file io_spinor.h.

◆ WRITE_SPINOR_SHARED_DOUBLE2

#define WRITE_SPINOR_SHARED_DOUBLE2   WRITE_SPINOR_SHARED_REAL

Definition at line 1009 of file io_spinor.h.

◆ WRITE_SPINOR_SHARED_FLOAT4

#define WRITE_SPINOR_SHARED_FLOAT4   WRITE_SPINOR_SHARED_REAL

Definition at line 1030 of file io_spinor.h.

◆ WRITE_SPINOR_SHARED_REAL

#define WRITE_SPINOR_SHARED_REAL (   tx,
  ty,
  tz,
  reg 
)
Value:
extern __shared__ char s_data[]; \
spinorFloat *sh = (spinorFloat*)s_data + DSLASH_SHARED_FLOATS_PER_THREAD*SHARED_STRIDE* \
((tx+blockDim.x*(ty+blockDim.y*tz))/SHARED_STRIDE) + ((tx+blockDim.x*(ty+blockDim.y*tz)) % SHARED_STRIDE); \
sh[0*SHARED_STRIDE] = reg##00_re; \
sh[1*SHARED_STRIDE] = reg##00_im; \
sh[2*SHARED_STRIDE] = reg##01_re; \
sh[3*SHARED_STRIDE] = reg##01_im; \
sh[4*SHARED_STRIDE] = reg##02_re; \
sh[5*SHARED_STRIDE] = reg##02_im; \
sh[6*SHARED_STRIDE] = reg##10_re; \
sh[7*SHARED_STRIDE] = reg##10_im; \
sh[8*SHARED_STRIDE] = reg##11_re; \
sh[9*SHARED_STRIDE] = reg##11_im; \
sh[10*SHARED_STRIDE] = reg##12_re; \
sh[11*SHARED_STRIDE] = reg##12_im; \
sh[12*SHARED_STRIDE] = reg##20_re; \
sh[13*SHARED_STRIDE] = reg##20_im; \
sh[14*SHARED_STRIDE] = reg##21_re; \
sh[15*SHARED_STRIDE] = reg##21_im; \
sh[16*SHARED_STRIDE] = reg##22_re; \
sh[17*SHARED_STRIDE] = reg##22_im; \
sh[18*SHARED_STRIDE] = reg##30_re; \
sh[19*SHARED_STRIDE] = reg##30_im; \
sh[20*SHARED_STRIDE] = reg##31_re; \
sh[21*SHARED_STRIDE] = reg##31_im; \
sh[22*SHARED_STRIDE] = reg##32_re; \
sh[23*SHARED_STRIDE] = reg##32_im;
blockDim
dim3 dim3 blockDim
Definition: CMakeCUDACompilerId.cpp1.ii:2471
tz
int tz
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1810
DSLASH_SHARED_FLOATS_PER_THREAD
#define DSLASH_SHARED_FLOATS_PER_THREAD
Definition: dslash_clover.cu:51
SHARED_STRIDE
#define SHARED_STRIDE
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:386
ty
int ty
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1366
spinorFloat
#define spinorFloat
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:65
tx
int tx
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:1146

Definition at line 980 of file io_spinor.h.

◆ WRITE_SPINOR_SHORT4

#define WRITE_SPINOR_SHORT4 (   stride)

Definition at line 377 of file io_spinor.h.

◆ WRITE_SPINOR_SHORT4_STR

#define WRITE_SPINOR_SHORT4_STR (   stride)

Definition at line 441 of file io_spinor.h.

◆ WRITE_ST_SPINOR_DOUBLE2

#define WRITE_ST_SPINOR_DOUBLE2 (   out,
  sid,
  mystride 
)
Value:
((double2*)out)[0*mystride+sid] = make_double2(o00_re, o00_im); \
((double2*)out)[1*mystride+sid] = make_double2(o01_re, o01_im); \
((double2*)out)[2*mystride+sid] = make_double2(o02_re, o02_im);
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359

Definition at line 843 of file io_spinor.h.

◆ WRITE_ST_SPINOR_DOUBLE2_STR

#define WRITE_ST_SPINOR_DOUBLE2_STR (   out,
  sid,
  mystride 
)
Value:
store_streaming_double2(&out[0*mystride+sid], o00_re, o00_im); \
store_streaming_double2(&out[1*mystride+sid], o01_re, o01_im); \
store_streaming_double2(&out[2*mystride+sid], o02_re, o02_im);
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
quda::store_streaming_double2
__device__ void store_streaming_double2(double2 *addr, double x, double y)
Definition: inline_ptx.h:49
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359

Definition at line 868 of file io_spinor.h.

◆ WRITE_ST_SPINOR_FLOAT2

#define WRITE_ST_SPINOR_FLOAT2 (   out,
  sid,
  mystride 
)
Value:
((float2*)out)[0*mystride+sid] = make_float2(o00_re, o00_im); \
((float2*)out)[1*mystride+sid] = make_float2(o01_re, o01_im); \
((float2*)out)[2*mystride+sid] = make_float2(o02_re, o02_im);
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359

Definition at line 848 of file io_spinor.h.

◆ WRITE_ST_SPINOR_FLOAT2_STR

#define WRITE_ST_SPINOR_FLOAT2_STR (   out,
  sid,
  mystride 
)
Value:
store_streaming_float2(&out[0*mystride+sid], o00_re, o00_im); \
store_streaming_float2(&out[1*mystride+sid], o01_re, o01_im); \
store_streaming_float2(&out[2*mystride+sid], o02_re, o02_im);
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
quda::store_streaming_float2
__device__ void store_streaming_float2(float2 *addr, float x, float y)
Definition: inline_ptx.h:54
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359

Definition at line 873 of file io_spinor.h.

◆ WRITE_ST_SPINOR_SHORT2

#define WRITE_ST_SPINOR_SHORT2 (   out,
  sid,
  mystride 
)
Value:
float c0 = fmaxf(fabsf(o00_re), fabsf(o00_im)); \
float c1 = fmaxf(fabsf(o01_re), fabsf(o01_im)); \
float c2 = fmaxf(fabsf(o02_re), fabsf(o02_im)); \
c0 = fmaxf(c0, c1); \
c0 = fmaxf(c0, c2); \
out ## Norm[sid] = c0; \
float scale = __fdividef(MAX_SHORT, c0); \
o00_re *= scale; o00_im *= scale; o01_re *= scale; o01_im *= scale; \
o02_re *= scale; o02_im *= scale; \
((short2*)out)[sid+0*mystride] = make_short2(f2i(o00_re), f2i(o00_im)); \
((short2*)out)[sid+1*mystride] = make_short2(f2i(o01_re), f2i(o01_im)); \
((short2*)out)[sid+2*mystride] = make_short2(f2i(o02_re), f2i(o02_im));
f2i
__device__ int f2i(float f)
Definition: convert.h:47
fabsf
float fabsf(float)
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
fmaxf
float fmaxf(float, float)
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
MAX_SHORT
#define MAX_SHORT
Definition: quda_internal.h:29
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359

Definition at line 853 of file io_spinor.h.

◆ WRITE_ST_SPINOR_SHORT2_STR

#define WRITE_ST_SPINOR_SHORT2_STR (   out,
  sid,
  mystride 
)
Value:
float c0 = fmaxf(fabsf(o00_re), fabsf(o00_im)); \
float c1 = fmaxf(fabsf(o01_re), fabsf(o01_im)); \
float c2 = fmaxf(fabsf(o02_re), fabsf(o02_im)); \
c0 = fmaxf(c0, c1); \
c0 = fmaxf(c0, c2); \
out ## Norm[sid] = c0; \
float scale = __fdividef(MAX_SHORT, c0); \
o00_re *= scale; o00_im *= scale; o01_re *= scale; o01_im *= scale; \
o02_re *= scale; o02_im *= scale; \
store_streaming_short2(&g_out[0*mystride+sid], f2i(o00_re), f2i(o00_im)); \
store_streaming_short2(&g_out[1*mystride+sid], f2i(o01_re), f2i(o01_im)); \
store_streaming_short2(&g_out[2*mystride+sid], f2i(o02_re), f2i(o02_im));
f2i
__device__ int f2i(float f)
Definition: convert.h:47
fabsf
float fabsf(float)
o01_re
VOLATILE spinorFloat o01_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:360
sid
int sid
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:396
o01_im
VOLATILE spinorFloat o01_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:361
o02_re
VOLATILE spinorFloat o02_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:362
o02_im
VOLATILE spinorFloat o02_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:363
o00_re
VOLATILE spinorFloat o00_re
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:358
fmaxf
float fmaxf(float, float)
MAX_SHORT
#define MAX_SHORT
Definition: quda_internal.h:29
o00_im
VOLATILE spinorFloat o00_im
Definition: asym_wilson_clover_dslash_dagger_fermi_core.h:359

Definition at line 878 of file io_spinor.h.

Generated by   doxygen 1.8.14