QUDA v0.4.0
A library for QCD on GPUs
Defines
quda/lib/io_spinor.h File Reference

Go to the source code of this file.

Defines

#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 READ_1ST_NBR_SPINOR_DOUBLE_TEX(spinor, idx, mystride)
#define READ_3RD_NBR_SPINOR_DOUBLE_TEX(spinor, idx, mystride)
#define READ_1ST_NBR_SPINOR_SINGLE_TEX(spinor, idx, mystride)
#define READ_3RD_NBR_SPINOR_SINGLE_TEX(spinor, idx, mystride)
#define READ_1ST_NBR_SPINOR_HALF_TEX(spinor, idx, mystride)
#define READ_3RD_NBR_SPINOR_HALF_TEX(spinor, idx, mystride)
#define READ_1ST_NBR_SPINOR_DOUBLE(spinor, idx, mystride)
#define READ_3RD_NBR_SPINOR_DOUBLE(spinor, idx, mystride)
#define READ_1ST_NBR_SPINOR_SINGLE(spinor, idx, mystride)
#define READ_3RD_NBR_SPINOR_SINGLE(spinor, idx, mystride)
#define READ_1ST_NBR_SPINOR_HALF(spinor, idx, mystride)
#define READ_3RD_NBR_SPINOR_HALF(spinor, idx, mystride)
#define WRITE_ST_SPINOR_DOUBLE2(out)
#define WRITE_ST_SPINOR_FLOAT2(out)
#define WRITE_ST_SPINOR_SHORT2(out)
#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)
#define READ_AND_SUM_ST_SPINOR_SINGLE_TEX(spinor)
#define READ_AND_SUM_ST_SPINOR_HALF_TEX_(spinor)
#define READ_AND_SUM_ST_SPINOR_HALF_TEX(spinor)   READ_AND_SUM_ST_SPINOR_HALF_TEX_(spinor)
#define READ_AND_SUM_ST_SPINOR(spinor)
#define READ_AND_SUM_ST_SPINOR_HALF_(spinor)
#define READ_AND_SUM_ST_SPINOR_HALF(spinor)   READ_AND_SUM_ST_SPINOR_HALF_(spinor)
#define READ_ST_ACCUM_DOUBLE_TEX(spinor)
#define READ_ST_ACCUM_SINGLE_TEX(spinor)
#define READ_ST_ACCUM_HALF_TEX(spinor)
#define READ_ST_ACCUM_DOUBLE(spinor)
#define READ_ST_ACCUM_SINGLE(spinor)
#define READ_ST_ACCUM_HALF(spinor)
#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 Documentation

#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];

Definition at line 469 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);

Definition at line 429 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);               \
  }

Definition at line 489 of file io_spinor.h.

#define READ_1ST_NBR_SPINOR_HALF_TEX (   spinor,
  idx,
  mystride 
)
Value:
float2 I0 = tex1Dfetch((spinor), idx + 0*mystride);                     \
  float2 I1 = tex1Dfetch((spinor), idx + 1*mystride);                   \
  float2 I2 = tex1Dfetch((spinor), idx + 2*mystride);                   \
  {                                                                     \
    float C = tex1Dfetch((spinorTexHalfNorm), norm_idx1);               \
    I0.x *= C; I0.y *= C;                                               \
    I1.x *= C; I1.y *= C;                                               \
    I2.x *= C; I2.y *= C;}

Definition at line 449 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];

Definition at line 479 of file io_spinor.h.

#define READ_1ST_NBR_SPINOR_SINGLE_TEX (   spinor,
  idx,
  mystride 
)
Value:
float2 I0 = tex1Dfetch((spinor), idx + 0*mystride);             \
  float2 I1 = tex1Dfetch((spinor), idx + 1*mystride);           \
  float2 I2 = tex1Dfetch((spinor), idx + 2*mystride);

Definition at line 439 of file io_spinor.h.

#define READ_3RD_NBR_SPINOR_DOUBLE (   spinor,
  idx,
  mystride 
)
Value:
double2 T0 = spinor[idx + 0*mystride];                  \
  double2 T1 = spinor[idx + 1*mystride];                        \
  double2 T2 = spinor[idx + 2*mystride];

Definition at line 474 of file io_spinor.h.

#define READ_3RD_NBR_SPINOR_DOUBLE_TEX (   spinor,
  idx,
  mystride 
)
Value:
double2 T0 = fetch_double2((spinor), idx + 0*mystride); \
  double2 T1 = fetch_double2((spinor), idx + 1*mystride);       \
  double2 T2 = fetch_double2((spinor), idx + 2*mystride);

Definition at line 434 of file io_spinor.h.

#define READ_3RD_NBR_SPINOR_HALF (   spinor,
  idx,
  mystride 
)
Value:
float2 T0, T1, T2;                                                      \
  {                                                                     \
    short2 S0 = in[idx + 0*mystride];                                   \
    short2 S1 = in[idx + 1*mystride];                                   \
    short2 S2 = in[idx + 2*mystride];                                   \
    float C = inNorm[idx];                                              \
    T0.x =C*short2float(S0.x); T0.y =C*short2float(S0.y);               \
    T1.x =C*short2float(S1.x); T1.y =C*short2float(S1.y);               \
    T2.x =C*short2float(S2.x); T2.y =C*short2float(S2.y);               \
  }

Definition at line 501 of file io_spinor.h.

#define READ_3RD_NBR_SPINOR_HALF_TEX (   spinor,
  idx,
  mystride 
)
Value:
float2 T0 = tex1Dfetch((spinor), idx + 0*mystride);                     \
  float2 T1 = tex1Dfetch((spinor), idx + 1*mystride);                   \
  float2 T2 = tex1Dfetch((spinor), idx + 2*mystride);                   \
  {                                                                     \
    float C = tex1Dfetch((spinorTexHalfNorm), norm_idx3);               \
    T0.x *= C; T0.y *= C;                                               \
    T1.x *= C; T1.y *= C;                                               \
    T2.x *= C; T2.y *= C;}

Definition at line 459 of file io_spinor.h.

#define READ_3RD_NBR_SPINOR_SINGLE (   spinor,
  idx,
  mystride 
)
Value:
float2 T0 = spinor[idx + 0*mystride];                           \
  float2 T1 = spinor[idx + 1*mystride];                         \
  float2 T2 = spinor[idx + 2*mystride];

Definition at line 484 of file io_spinor.h.

#define READ_3RD_NBR_SPINOR_SINGLE_TEX (   spinor,
  idx,
  mystride 
)
Value:
float2 T0 = tex1Dfetch((spinor), idx + 0*mystride);             \
  float2 T1 = tex1Dfetch((spinor), idx + 1*mystride);           \
  float2 T2 = tex1Dfetch((spinor), idx + 2*mystride);

Definition at line 444 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];

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));

Definition at line 160 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;

Definition at line 113 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((spinor), sid + 0*(stride));    \
  float4 accum1 = tex1Dfetch((spinor), sid + 1*(stride));          \
  float4 accum2 = tex1Dfetch((spinor), sid + 2*(stride));          \
  float4 accum3 = tex1Dfetch((spinor), sid + 3*(stride));          \
  float4 accum4 = tex1Dfetch((spinor), sid + 4*(stride));          \
  float4 accum5 = tex1Dfetch((spinor), sid + 5*(stride));          \
  float C = tex1Dfetch((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;

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)];

Definition at line 105 of file io_spinor.h.

#define READ_ACCUM_SINGLE_TEX (   spinor,
  stride 
)
Value:
float4 accum0 = tex1Dfetch((spinor), sid + 0*(stride));       \
  float4 accum1 = tex1Dfetch((spinor), sid + 1*(stride));       \
  float4 accum2 = tex1Dfetch((spinor), sid + 2*(stride));       \
  float4 accum3 = tex1Dfetch((spinor), sid + 3*(stride));       \
  float4 accum4 = tex1Dfetch((spinor), sid + 4*(stride));       \
  float4 accum5 = tex1Dfetch((spinor), sid + 5*(stride));

Definition at line 234 of file io_spinor.h.

#define READ_AND_SUM_ST_SPINOR (   spinor)
Value:
o00_re += spinor[0*sp_stride+sid].x; o00_im += spinor[0*sp_stride+sid].y; \
  o01_re += spinor[1*sp_stride+sid].x; o01_im += spinor[1*sp_stride+sid].y; \
  o02_re += spinor[2*sp_stride+sid].x; o02_im += spinor[2*sp_stride+sid].y; \

Definition at line 600 of file io_spinor.h.

#define READ_AND_SUM_ST_SPINOR_DOUBLE_TEX (   spinor)
Value:
{                       \
  double2 tmp0 = fetch_double2((spinor), sid + 0*(sp_stride));          \
  double2 tmp1 = fetch_double2((spinor), sid + 1*(sp_stride));          \
  double2 tmp2 = fetch_double2((spinor), sid + 2*(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; }

Definition at line 572 of file io_spinor.h.

#define READ_AND_SUM_ST_SPINOR_HALF (   spinor)    READ_AND_SUM_ST_SPINOR_HALF_(spinor)

Definition at line 614 of file io_spinor.h.

#define READ_AND_SUM_ST_SPINOR_HALF_ (   spinor)
Value:
float C = spinor ## Norm[sid];                          \
  o00_re += C*short2float(spinor[0*sp_stride + sid].x);         \
  o00_im += C*short2float(spinor[0*sp_stride + sid].y);         \
  o01_re += C*short2float(spinor[1*sp_stride + sid].x);         \
  o01_im += C*short2float(spinor[1*sp_stride + sid].y);         \
  o02_re += C*short2float(spinor[2*sp_stride + sid].x);         \
  o02_im += C*short2float(spinor[2*sp_stride + sid].y);

Definition at line 605 of file io_spinor.h.

#define READ_AND_SUM_ST_SPINOR_HALF_TEX (   spinor)    READ_AND_SUM_ST_SPINOR_HALF_TEX_(spinor)

Definition at line 597 of file io_spinor.h.

#define READ_AND_SUM_ST_SPINOR_HALF_TEX_ (   spinor)
Value:
{                       \
  float2 tmp0 = tex1Dfetch((spinor), sid + 0*sp_stride);                \
  float2 tmp1 = tex1Dfetch((spinor), sid + 1*sp_stride);                \
  float2 tmp2 = tex1Dfetch((spinor), sid + 2*sp_stride);                \
  float C = tex1Dfetch((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; }

Definition at line 588 of file io_spinor.h.

#define READ_AND_SUM_ST_SPINOR_SINGLE_TEX (   spinor)
Value:
{                       \
  float2 tmp0 = tex1Dfetch((spinor), sid + 0*(sp_stride));              \
  float2 tmp1 = tex1Dfetch((spinor), sid + 1*(sp_stride));              \
  float2 tmp2 = tex1Dfetch((spinor), sid + 2*(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; }

Definition at line 580 of file io_spinor.h.

#define READ_HALF_SPINOR   READ_SPINOR_UP

Definition at line 390 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)];

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)];

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));

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));

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)];

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));

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;

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;

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((spinor), sp_idx + 3*(stride));      \
  float4 I4 = tex1Dfetch((spinor), sp_idx + 4*(stride));      \
  float4 I5 = tex1Dfetch((spinor), sp_idx + 5*(stride));      \
  float C = tex1Dfetch((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;

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((spinor), sp_idx + 0*(stride));   \
  float4 I1 = tex1Dfetch((spinor), sp_idx + 1*(stride));   \
  float4 I2 = tex1Dfetch((spinor), sp_idx + 2*(stride));   \
  float4 I3 = tex1Dfetch((spinor), sp_idx + 3*(stride));   \
  float4 I4 = tex1Dfetch((spinor), sp_idx + 4*(stride));   \
  float4 I5 = tex1Dfetch((spinor), sp_idx + 5*(stride));   \
  float C = tex1Dfetch((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;

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;                         \

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((spinor), sp_idx + 0*(stride));      \
  float4 I1 = tex1Dfetch((spinor), sp_idx + 1*(stride));      \
  float4 I2 = tex1Dfetch((spinor), sp_idx + 2*(stride));      \
  float C = tex1Dfetch((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;                 \

Definition at line 210 of file io_spinor.h.

#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]);

Definition at line 689 of file io_spinor.h.

#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]);

Definition at line 710 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)];

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)];

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((spinor), sp_idx + 3*(stride));   \
  float4 I4 = tex1Dfetch((spinor), sp_idx + 4*(stride));   \
  float4 I5 = tex1Dfetch((spinor), sp_idx + 5*(stride));

Definition at line 187 of file io_spinor.h.

#define READ_SPINOR_SINGLE_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 I0 = tex1Dfetch((spinor), sp_idx + 0*(stride));   \
  float4 I1 = tex1Dfetch((spinor), sp_idx + 1*(stride));   \
  float4 I2 = tex1Dfetch((spinor), sp_idx + 2*(stride));   \
  float4 I3 = tex1Dfetch((spinor), sp_idx + 3*(stride));   \
  float4 I4 = tex1Dfetch((spinor), sp_idx + 4*(stride));   \
  float4 I5 = tex1Dfetch((spinor), sp_idx + 5*(stride));

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)];   \

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((spinor), sp_idx + 0*(stride));   \
  float4 I1 = tex1Dfetch((spinor), sp_idx + 1*(stride));   \
  float4 I2 = tex1Dfetch((spinor), sp_idx + 2*(stride));   \

Definition at line 182 of file io_spinor.h.

#define READ_ST_ACCUM_DOUBLE (   spinor)
Value:
double2 accum0 = spinor[sid + 0*(sp_stride)];                      \
  double2 accum1 = spinor[sid + 1*(sp_stride)];                    \
  double2 accum2 = spinor[sid + 2*(sp_stride)];

Definition at line 636 of file io_spinor.h.

#define READ_ST_ACCUM_DOUBLE_TEX (   spinor)
Value:
double2 accum0 = fetch_double2((spinor), sid + 0*(sp_stride));   \
  double2 accum1 = fetch_double2((spinor), sid + 1*(sp_stride));   \
  double2 accum2 = fetch_double2((spinor), sid + 2*(sp_stride));

Definition at line 617 of file io_spinor.h.

#define READ_ST_ACCUM_HALF (   spinor)
Value:
float2 accum0, accum1, accum2;                                  \
  {                                                                     \
    short2 S0 = x[sid + 0*sp_stride];                                   \
    short2 S1 = x[sid + 1*sp_stride];                                   \
    short2 S2 = x[sid + 2*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);       \
  }

Definition at line 646 of file io_spinor.h.

#define READ_ST_ACCUM_HALF_TEX (   spinor)
Value:
float2 accum0 = tex1Dfetch((spinor), sid + 0*sp_stride);        \
  float2 accum1 = tex1Dfetch((spinor), sid + 1*sp_stride);      \
  float2 accum2 = tex1Dfetch((spinor), sid + 2*sp_stride);      \
  float C = tex1Dfetch((accumTexHalfNorm), sid);                \
  accum0.x *= C; accum0.y *= C;                                 \
  accum1.x *= C; accum1.y *= C;                                 \
  accum2.x *= C; accum2.y *= C;

Definition at line 627 of file io_spinor.h.

#define READ_ST_ACCUM_SINGLE (   spinor)
Value:
float2 accum0 = spinor[sid + 0*(sp_stride)];                    \
  float2 accum1 = spinor[sid + 1*(sp_stride)];                  \
  float2 accum2 = spinor[sid + 2*(sp_stride)];

Definition at line 641 of file io_spinor.h.

#define READ_ST_ACCUM_SINGLE_TEX (   spinor)
Value:
float2 accum0 = tex1Dfetch((spinor), sid + 0*sp_stride);        \
  float2 accum1 = tex1Dfetch((spinor), sid + 1*sp_stride);      \
  float2 accum2 = tex1Dfetch((spinor), sid + 2*sp_stride);

Definition at line 622 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);

Definition at line 392 of file io_spinor.h.

#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);

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);

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);

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);

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

#define WRITE_SPINOR_SHARED_FLOAT4   WRITE_SPINOR_SHARED_REAL

Definition at line 708 of file io_spinor.h.

#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;

Definition at line 658 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)
Value:
out[0*sp_stride+sid] = make_double2(o00_re, o00_im);    \
  out[1*sp_stride+sid] = make_double2(o01_re, o01_im);  \
  out[2*sp_stride+sid] = make_double2(o02_re, o02_im);

Definition at line 514 of file io_spinor.h.

#define WRITE_ST_SPINOR_DOUBLE2_STR ( )    WRITE_ST_SPINOR_DOUBLE2()

Definition at line 566 of file io_spinor.h.

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

Definition at line 519 of file io_spinor.h.

#define WRITE_ST_SPINOR_FLOAT4_STR ( )    WRITE_ST_SPINOR_FLOAT4()

Definition at line 567 of file io_spinor.h.

#define WRITE_ST_SPINOR_SHORT2 (   out)
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*sp_stride] = make_short2((short)o00_re, (short)o00_im);     \
  out[sid+1*sp_stride] = make_short2((short)o01_re, (short)o01_im);     \
  out[sid+2*sp_stride] = make_short2((short)o02_re, (short)o02_im);

Definition at line 524 of file io_spinor.h.

#define WRITE_ST_SPINOR_SHORT4_STR ( )    WRITE_ST_SPINOR_SHORT4()

Definition at line 568 of file io_spinor.h.

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