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

Go to the source code of this file.

Macros

#define READ_SPINOR_DOUBLE(spinor, stride, sp_idx, norm_idx)
 
#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_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_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_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_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_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)   WRITE_SPINOR_DOUBLE2(stride)
 
#define WRITE_SPINOR_FLOAT4_STR(stride)   WRITE_SPINOR_FLOAT4(stride)
 
#define WRITE_SPINOR_SHORT4_STR(stride)   WRITE_SPINOR_SHORT4(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_KS_NBR_SPINOR_DOUBLE_TEX(T, spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_SINGLE_TEX(spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_SINGLE_TEX(T, spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_HALF_TEX_(spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_HALF_TEX(spinor, idx, mystride)   READ_1ST_NBR_SPINOR_HALF_TEX_(spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_HALF_TEX_(T, spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_HALF_TEX(T, spinor, idx, mystride)   READ_KS_NBR_SPINOR_HALF_TEX_(T, spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_DOUBLE(spinor, idx, mystride)
 
#define READ_KS_NBR_SPINOR_DOUBLE(T, spinor, idx, mystride)
 
#define READ_1ST_NBR_SPINOR_SINGLE(spinor, idx, mystride)
 
#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()   WRITE_ST_SPINOR_DOUBLE2()
 
#define WRITE_ST_SPINOR_FLOAT4_STR()   WRITE_ST_SPINOR_FLOAT4()
 
#define WRITE_ST_SPINOR_SHORT4_STR()   WRITE_ST_SPINOR_SHORT4()
 
#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

#define APPLY_TWIST (   a,
  reg 
)

Definition at line 1187 of file io_spinor.h.

#define APPLY_TWIST_INV (   a,
  b,
  reg 
)

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

Definition at line 1122 of file io_spinor.h.

#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];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 1033 of file io_spinor.h.

#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));
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 1077 of file io_spinor.h.

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

Definition at line 1073 of file io_spinor.h.

#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; \
}
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 1055 of file io_spinor.h.

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

Definition at line 1117 of file io_spinor.h.

#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; \
}
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 1100 of file io_spinor.h.

#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)];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 1047 of file io_spinor.h.

#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));
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 1092 of file io_spinor.h.

#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];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 598 of file io_spinor.h.

#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);
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 549 of file io_spinor.h.

#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); \
}
cpuColorSpinorField * in

Definition at line 618 of file io_spinor.h.

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

Definition at line 582 of file io_spinor.h.

#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;}
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 572 of file io_spinor.h.

#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];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 608 of file io_spinor.h.

#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);
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 559 of file io_spinor.h.

#define READ_ACCUM_DOUBLE (   spinor,
  stride 
)
Value:
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];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 91 of file io_spinor.h.

#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));
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 160 of file io_spinor.h.

#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];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

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

Definition at line 886 of file io_spinor.h.

#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));
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 959 of file io_spinor.h.

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

Definition at line 956 of file io_spinor.h.

#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;
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 928 of file io_spinor.h.

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

Definition at line 1029 of file io_spinor.h.

#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)];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 913 of file io_spinor.h.

#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));
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 986 of file io_spinor.h.

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

Definition at line 128 of file io_spinor.h.

#define READ_ACCUM_HALF_ (   spinor,
  stride 
)
Value:
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;
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 113 of file io_spinor.h.

#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;
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 1000 of file io_spinor.h.

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

Definition at line 257 of file io_spinor.h.

#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;
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 242 of file io_spinor.h.

#define READ_ACCUM_SINGLE (   spinor,
  stride 
)
Value:
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)];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 105 of file io_spinor.h.

#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));
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 234 of file io_spinor.h.

#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; \
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17

Definition at line 728 of file io_spinor.h.

#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; }
cudaColorSpinorField * tmp1
Definition: dslash_test.cpp:41
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17
cudaColorSpinorField * tmp2
Definition: dslash_test.cpp:41
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 700 of file io_spinor.h.

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

Definition at line 742 of file io_spinor.h.

#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);
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17

Definition at line 733 of file io_spinor.h.

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

Definition at line 725 of file io_spinor.h.

#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; }
cudaColorSpinorField * tmp1
Definition: dslash_test.cpp:41
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17
cudaColorSpinorField * tmp2
Definition: dslash_test.cpp:41
#define TEX1DFETCH(type, tex, idx)

Definition at line 716 of file io_spinor.h.

#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; }
cudaColorSpinorField * tmp1
Definition: dslash_test.cpp:41
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17
cudaColorSpinorField * tmp2
Definition: dslash_test.cpp:41
#define TEX1DFETCH(type, tex, idx)

Definition at line 708 of file io_spinor.h.

#define READ_HALF_SPINOR   READ_SPINOR_UP

Definition at line 390 of file io_spinor.h.

#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];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 603 of file io_spinor.h.

#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);
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 554 of file io_spinor.h.

#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); \
}
int y[4]
cpuColorSpinorField * in
int x[4]

Definition at line 630 of file io_spinor.h.

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

Definition at line 595 of file io_spinor.h.

#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;}
int y[4]
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
int x[4]
#define TEX1DFETCH(type, tex, idx)

Definition at line 585 of file io_spinor.h.

#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];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 613 of file io_spinor.h.

#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);
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 564 of file io_spinor.h.

#define READ_SPINOR_DOUBLE (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
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)];
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 1 of file io_spinor.h.

#define READ_SPINOR_DOUBLE_DOWN (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
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)];
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 23 of file io_spinor.h.

#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));
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 152 of file io_spinor.h.

#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));
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 130 of file io_spinor.h.

#define READ_SPINOR_DOUBLE_UP (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
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)];
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 15 of file io_spinor.h.

#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));
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 144 of file io_spinor.h.

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

Definition at line 64 of file io_spinor.h.

#define READ_SPINOR_HALF_ (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
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;
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 49 of file io_spinor.h.

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

Definition at line 88 of file io_spinor.h.

#define READ_SPINOR_HALF_DOWN_ (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
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;
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 79 of file io_spinor.h.

#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 231 of file io_spinor.h.

#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;
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 222 of file io_spinor.h.

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

Definition at line 207 of file io_spinor.h.

#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;
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 192 of file io_spinor.h.

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

Definition at line 76 of file io_spinor.h.

#define READ_SPINOR_HALF_UP_ (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
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; \
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 67 of file io_spinor.h.

#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 219 of file io_spinor.h.

#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; \
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 210 of file io_spinor.h.

#define READ_SPINOR_SHARED_DOUBLE2 (   tx,
  ty,
  tz 
)
Value:
extern __shared__ char s_data[]; \
((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]);
#define DSLASH_SHARED_FLOATS_PER_THREAD
__shared__ char s_data[]

Definition at line 819 of file io_spinor.h.

#define READ_SPINOR_SHARED_FLOAT4 (   tx,
  ty,
  tz 
)
Value:
extern __shared__ char s_data[]; \
((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]);
#define DSLASH_SHARED_FLOATS_PER_THREAD
__shared__ char s_data[]

Definition at line 840 of file io_spinor.h.

#define READ_SPINOR_SINGLE (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
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)];
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 31 of file io_spinor.h.

#define READ_SPINOR_SINGLE_DOWN (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I3 = spinor[sp_idx + 3*(stride)]; \
float4 I4 = spinor[sp_idx + 4*(stride)]; \
float4 I5 = spinor[sp_idx + 5*(stride)];
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 44 of file io_spinor.h.

#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));
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 187 of file io_spinor.h.

#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));
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 174 of file io_spinor.h.

#define READ_SPINOR_SINGLE_UP (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I0 = spinor[sp_idx + 0*(stride)]; \
float4 I1 = spinor[sp_idx + 1*(stride)]; \
float4 I2 = spinor[sp_idx + 2*(stride)]; \
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40

Definition at line 39 of file io_spinor.h.

#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)); \
int sp_idx
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
#define TEX1DFETCH(type, tex, idx)

Definition at line 182 of file io_spinor.h.

#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)];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17

Definition at line 766 of file io_spinor.h.

#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));
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:90

Definition at line 745 of file io_spinor.h.

#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); \
}
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17
int x[4]

Definition at line 776 of file io_spinor.h.

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

Definition at line 764 of file io_spinor.h.

#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;
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17
#define TEX1DFETCH(type, tex, idx)

Definition at line 755 of file io_spinor.h.

#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)];
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17

Definition at line 771 of file io_spinor.h.

#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);
cpuColorSpinorField * spinor
Definition: dslash_test.cpp:40
QudaGaugeParam param
Definition: pack_test.cpp:17
#define TEX1DFETCH(type, tex, idx)

Definition at line 750 of file io_spinor.h.

#define WRITE_FLAVOR_SPINOR_DOUBLE2 ( )
Value:
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);
VOLATILE spinorFloat o2_01_re
VOLATILE spinorFloat o1_30_re
VOLATILE spinorFloat o1_02_im
VOLATILE spinorFloat o1_32_re
VOLATILE spinorFloat o1_22_im
VOLATILE spinorFloat o2_32_im
VOLATILE spinorFloat o2_02_im
VOLATILE spinorFloat o2_31_re
VOLATILE spinorFloat o1_31_re
VOLATILE spinorFloat o2_32_re
VOLATILE spinorFloat o1_20_re
VOLATILE spinorFloat o1_02_re
VOLATILE spinorFloat o2_22_re
VOLATILE spinorFloat o1_10_im
VOLATILE spinorFloat o2_30_re
VOLATILE spinorFloat o2_31_im
VOLATILE spinorFloat o1_31_im
QudaGaugeParam param
Definition: pack_test.cpp:17
VOLATILE spinorFloat o1_22_re
VOLATILE spinorFloat o1_00_re
VOLATILE spinorFloat o1_12_re
VOLATILE spinorFloat o2_12_re
VOLATILE spinorFloat o1_11_re
VOLATILE spinorFloat o2_22_im
VOLATILE spinorFloat o2_00_im
VOLATILE spinorFloat o2_21_im
VOLATILE spinorFloat o1_32_im
VOLATILE spinorFloat o1_21_im
VOLATILE spinorFloat o2_10_im
VOLATILE spinorFloat o1_01_re
VOLATILE spinorFloat o2_01_im
VOLATILE spinorFloat o1_01_im
VOLATILE spinorFloat o1_11_im
VOLATILE spinorFloat o1_12_im
VOLATILE spinorFloat o2_10_re
VOLATILE spinorFloat o2_02_re
VOLATILE spinorFloat o1_20_im
VOLATILE spinorFloat o2_00_re
VOLATILE spinorFloat o2_11_im
VOLATILE spinorFloat o1_10_re
cpuColorSpinorField * out
VOLATILE spinorFloat o1_30_im
VOLATILE spinorFloat o2_20_im
VOLATILE spinorFloat o2_30_im
VOLATILE spinorFloat o2_12_im
VOLATILE spinorFloat o2_21_re
VOLATILE spinorFloat o1_21_re
VOLATILE spinorFloat o2_20_re
VOLATILE spinorFloat o2_11_re
VOLATILE spinorFloat o1_00_im

ndeg tm:

Definition at line 428 of file io_spinor.h.

#define WRITE_FLAVOR_SPINOR_FLOAT4 ( )
Value:
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);
VOLATILE spinorFloat o2_01_re
VOLATILE spinorFloat o1_30_re
VOLATILE spinorFloat o1_02_im
VOLATILE spinorFloat o1_32_re
VOLATILE spinorFloat o1_22_im
VOLATILE spinorFloat o2_32_im
VOLATILE spinorFloat o2_02_im
VOLATILE spinorFloat o2_31_re
VOLATILE spinorFloat o1_31_re
VOLATILE spinorFloat o2_32_re
VOLATILE spinorFloat o1_20_re
VOLATILE spinorFloat o1_02_re
VOLATILE spinorFloat o2_22_re
VOLATILE spinorFloat o1_10_im
VOLATILE spinorFloat o2_30_re
VOLATILE spinorFloat o2_31_im
VOLATILE spinorFloat o1_31_im
QudaGaugeParam param
Definition: pack_test.cpp:17
VOLATILE spinorFloat o1_22_re
VOLATILE spinorFloat o1_00_re
VOLATILE spinorFloat o1_12_re
VOLATILE spinorFloat o2_12_re
VOLATILE spinorFloat o1_11_re
VOLATILE spinorFloat o2_22_im
VOLATILE spinorFloat o2_00_im
VOLATILE spinorFloat o2_21_im
VOLATILE spinorFloat o1_32_im
VOLATILE spinorFloat o1_21_im
VOLATILE spinorFloat o2_10_im
VOLATILE spinorFloat o1_01_re
VOLATILE spinorFloat o2_01_im
VOLATILE spinorFloat o1_01_im
VOLATILE spinorFloat o1_11_im
VOLATILE spinorFloat o1_12_im
VOLATILE spinorFloat o2_10_re
VOLATILE spinorFloat o2_02_re
VOLATILE spinorFloat o1_20_im
VOLATILE spinorFloat o2_00_re
VOLATILE spinorFloat o2_11_im
VOLATILE spinorFloat o1_10_re
cpuColorSpinorField * out
VOLATILE spinorFloat o1_30_im
VOLATILE spinorFloat o2_20_im
VOLATILE spinorFloat o2_30_im
VOLATILE spinorFloat o2_12_im
VOLATILE spinorFloat o2_21_re
VOLATILE spinorFloat o1_21_re
VOLATILE spinorFloat o2_20_re
VOLATILE spinorFloat o2_11_re
VOLATILE spinorFloat o1_00_im

Definition at line 455 of file io_spinor.h.

#define WRITE_FLAVOR_SPINOR_SHORT4 ( )

Definition at line 470 of file io_spinor.h.

#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);
cpuColorSpinorField * out

Definition at line 392 of file io_spinor.h.

#define WRITE_HALF_SPINOR_FLOAT4 (   stride,
  sid 
)
Value:

Definition at line 400 of file io_spinor.h.

#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((short)a0_re, (short)a0_im, (short)a1_re, (short)a1_im); \
out[sid+1*(stride)] = make_short4((short)a2_re, (short)a2_im, (short)b0_re, (short)b0_im); \
out[sid+2*(stride)] = make_short4((short)b1_re, (short)b1_im, (short)b2_re, (short)b2_im);
cpuColorSpinorField * out
#define MAX_SHORT
Definition: quda_internal.h:30

Definition at line 405 of file io_spinor.h.

#define WRITE_SPINOR_DOUBLE2 (   stride)
Value:
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);
#define o32_im
Definition: gamma5.h:295
#define o31_im
Definition: gamma5.h:293
#define o30_im
Definition: gamma5.h:291
#define o32_re
Definition: gamma5.h:294
#define o31_re
Definition: gamma5.h:292
cpuColorSpinorField * out

Definition at line 260 of file io_spinor.h.

#define WRITE_SPINOR_DOUBLE2_STR (   stride)    WRITE_SPINOR_DOUBLE2(stride)

Definition at line 383 of file io_spinor.h.

#define WRITE_SPINOR_FLOAT4 (   stride)
Value:
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);
#define o32_im
Definition: gamma5.h:295
#define o31_im
Definition: gamma5.h:293
#define o30_im
Definition: gamma5.h:291
#define o32_re
Definition: gamma5.h:294
#define o31_re
Definition: gamma5.h:292
cpuColorSpinorField * out

Definition at line 274 of file io_spinor.h.

#define WRITE_SPINOR_FLOAT4_STR (   stride)    WRITE_SPINOR_FLOAT4(stride)

Definition at line 384 of file io_spinor.h.

#define WRITE_SPINOR_SHARED_DOUBLE2   WRITE_SPINOR_SHARED_REAL

Definition at line 817 of file io_spinor.h.

#define WRITE_SPINOR_SHARED_FLOAT4   WRITE_SPINOR_SHARED_REAL

Definition at line 838 of file io_spinor.h.

#define WRITE_SPINOR_SHARED_REAL (   tx,
  ty,
  tz,
  reg 
)
Value:
extern __shared__ char s_data[]; \
((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;
#define DSLASH_SHARED_FLOATS_PER_THREAD
__shared__ char s_data[]

Definition at line 788 of file io_spinor.h.

#define WRITE_SPINOR_SHORT4 (   stride)

Definition at line 282 of file io_spinor.h.

#define WRITE_SPINOR_SHORT4_STR (   stride)    WRITE_SPINOR_SHORT4(stride)

Definition at line 385 of file io_spinor.h.

#define WRITE_ST_SPINOR_DOUBLE2 (   out,
  sid,
  mystride 
)
Value:
out[0*mystride+sid] = make_double2(o00_re, o00_im); \
out[1*mystride+sid] = make_double2(o01_re, o01_im); \
out[2*mystride+sid] = make_double2(o02_re, o02_im);
cpuColorSpinorField * out

Definition at line 642 of file io_spinor.h.

#define WRITE_ST_SPINOR_DOUBLE2_STR ( )    WRITE_ST_SPINOR_DOUBLE2()

Definition at line 694 of file io_spinor.h.

#define WRITE_ST_SPINOR_FLOAT2 (   out,
  sid,
  mystride 
)
Value:
out[0*mystride+sid] = make_float2(o00_re, o00_im); \
out[1*mystride+sid] = make_float2(o01_re, o01_im); \
out[2*mystride+sid] = make_float2(o02_re, o02_im);
cpuColorSpinorField * out

Definition at line 647 of file io_spinor.h.

#define WRITE_ST_SPINOR_FLOAT4_STR ( )    WRITE_ST_SPINOR_FLOAT4()

Definition at line 695 of file io_spinor.h.

#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; \
out[sid+0*mystride] = make_short2((short)o00_re, (short)o00_im); \
out[sid+1*mystride] = make_short2((short)o01_re, (short)o01_im); \
out[sid+2*mystride] = make_short2((short)o02_re, (short)o02_im);
cpuColorSpinorField * out
#define MAX_SHORT
Definition: quda_internal.h:30

Definition at line 652 of file io_spinor.h.

#define WRITE_ST_SPINOR_SHORT4_STR ( )    WRITE_ST_SPINOR_SHORT4()

Definition at line 696 of file io_spinor.h.