QUDA  0.9.0
Macros | Functions
contract_core.h File Reference

Go to the source code of this file.

Macros

#define _TWIST_QUDA_CONTRACT
 
#define tmp_re   tmp.x
 
#define tmp_im   tmp.y
 
#define TOTAL_COMPONENTS   16
 
#define READ_INTERMEDIATE_SPINOR_DOUBLE(spinor, stride, sp_idx, norm_idx)
 
#define READ_INTERMEDIATE_SPINOR_DOUBLE_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR   READ_SPINOR_DOUBLE_TEX
 
#define READ_INTERMEDIATE_SPINOR   READ_INTERMEDIATE_SPINOR_DOUBLE_TEX
 
#define SPINORTEX   spinorTexDouble
 
#define INTERTEX   interTexDouble
 
#define SPINOR_HOP   12
 
#define READ_SPINOR_SINGLE(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR_SINGLE_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_INTERMEDIATE_SPINOR_SINGLE(spinor, stride, sp_idx, norm_idx)
 
#define READ_INTERMEDIATE_SPINOR_SINGLE_TEX(spinor, stride, sp_idx, norm_idx)
 
#define READ_SPINOR   READ_SPINOR_SINGLE_TEX
 
#define SPINORTEX   spinorTexSingle
 
#define READ_INTERMEDIATE_SPINOR   READ_INTERMEDIATE_SPINOR_SINGLE_TEX
 
#define INTERTEX   interTexSingle
 
#define SPINOR_HOP   6
 

Functions

__global__ void contractGamma5Kernel (double2 *out, double2 *in1, double2 *in2, int myStride, const int Parity, const DslashParam param)
 
__global__ void contractTsliceKernel (double2 *out, double2 *in1, double2 *in2, int myStride, const int Tslice, const int Parity, const DslashParam param)
 
__global__ void contractKernel (double2 *out, double2 *in1, double2 *in2, int myStride, const int Parity, const DslashParam param)
 
__global__ void contractGamma5Kernel (float2 *out, float4 *in1, float4 *in2, int myStride, const int Parity, const DslashParam param)
 
__global__ void contractTsliceKernel (float2 *out, float4 *in1, float4 *in2, int myStride, const int Tslice, const int Parity, const DslashParam param)
 
__global__ void contractKernel (float2 *out, float4 *in1, float4 *in2, int myStride, const int Parity, const DslashParam param)
 

Macro Definition Documentation

◆ _TWIST_QUDA_CONTRACT

#define _TWIST_QUDA_CONTRACT

Definition at line 6 of file contract_core.h.

◆ INTERTEX [1/2]

#define INTERTEX   interTexDouble

Definition at line 904 of file contract_core.h.

Referenced by contractGamma5Kernel(), contractKernel(), and contractTsliceKernel().

◆ INTERTEX [2/2]

#define INTERTEX   interTexSingle

Definition at line 904 of file contract_core.h.

◆ READ_INTERMEDIATE_SPINOR [1/2]

#define READ_INTERMEDIATE_SPINOR   READ_INTERMEDIATE_SPINOR_DOUBLE_TEX

Definition at line 899 of file contract_core.h.

Referenced by contractGamma5Kernel(), contractKernel(), and contractTsliceKernel().

◆ READ_INTERMEDIATE_SPINOR [2/2]

#define READ_INTERMEDIATE_SPINOR   READ_INTERMEDIATE_SPINOR_SINGLE_TEX

Definition at line 899 of file contract_core.h.

◆ READ_INTERMEDIATE_SPINOR_DOUBLE

#define READ_INTERMEDIATE_SPINOR_DOUBLE (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
double2 J0 = spinor[sp_idx + 0*(stride)]; \
double2 J1 = spinor[sp_idx + 1*(stride)]; \
double2 J2 = spinor[sp_idx + 2*(stride)]; \
double2 J3 = spinor[sp_idx + 3*(stride)]; \
double2 J4 = spinor[sp_idx + 4*(stride)]; \
double2 J5 = spinor[sp_idx + 5*(stride)]; \
double2 J6 = spinor[sp_idx + 6*(stride)]; \
double2 J7 = spinor[sp_idx + 7*(stride)]; \
double2 J8 = spinor[sp_idx + 8*(stride)]; \
double2 J9 = spinor[sp_idx + 9*(stride)]; \
double2 J10 = spinor[sp_idx +10*(stride)]; \
double2 J11 = spinor[sp_idx +11*(stride)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 13 of file contract_core.h.

◆ READ_INTERMEDIATE_SPINOR_DOUBLE_TEX

#define READ_INTERMEDIATE_SPINOR_DOUBLE_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
double2 J0 = fetch_double2((spinor), sp_idx + 0*(stride)); \
double2 J1 = fetch_double2((spinor), sp_idx + 1*(stride)); \
double2 J2 = fetch_double2((spinor), sp_idx + 2*(stride)); \
double2 J3 = fetch_double2((spinor), sp_idx + 3*(stride)); \
double2 J4 = fetch_double2((spinor), sp_idx + 4*(stride)); \
double2 J5 = fetch_double2((spinor), sp_idx + 5*(stride)); \
double2 J6 = fetch_double2((spinor), sp_idx + 6*(stride)); \
double2 J7 = fetch_double2((spinor), sp_idx + 7*(stride)); \
double2 J8 = fetch_double2((spinor), sp_idx + 8*(stride)); \
double2 J9 = fetch_double2((spinor), sp_idx + 9*(stride)); \
double2 J10 = fetch_double2((spinor), sp_idx +10*(stride)); \
double2 J11 = fetch_double2((spinor), sp_idx +11*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
fetch_double2
__inline__ __device__ double2 fetch_double2(texture< int4, 1 > t, int i)
Definition: texture.h:88
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 27 of file contract_core.h.

◆ READ_INTERMEDIATE_SPINOR_SINGLE

#define READ_INTERMEDIATE_SPINOR_SINGLE (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 J0 = spinor[sp_idx + 0*(stride)]; \
float4 J1 = spinor[sp_idx + 1*(stride)]; \
float4 J2 = spinor[sp_idx + 2*(stride)]; \
float4 J3 = spinor[sp_idx + 3*(stride)]; \
float4 J4 = spinor[sp_idx + 4*(stride)]; \
float4 J5 = spinor[sp_idx + 5*(stride)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 865 of file contract_core.h.

◆ READ_INTERMEDIATE_SPINOR_SINGLE_TEX

#define READ_INTERMEDIATE_SPINOR_SINGLE_TEX (   spinor,
  stride,
  sp_idx,
  norm_idx 
)
Value:
float4 J0 = TEX1DFETCH(float4, (spinor), sp_idx + 0*(stride)); \
float4 J1 = TEX1DFETCH(float4, (spinor), sp_idx + 1*(stride)); \
float4 J2 = TEX1DFETCH(float4, (spinor), sp_idx + 2*(stride)); \
float4 J3 = TEX1DFETCH(float4, (spinor), sp_idx + 3*(stride)); \
float4 J4 = TEX1DFETCH(float4, (spinor), sp_idx + 4*(stride)); \
float4 J5 = TEX1DFETCH(float4, (spinor), sp_idx + 5*(stride));
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41
TEX1DFETCH
#define TEX1DFETCH(type, tex, idx)
Definition: dslash_textures.h:7

Definition at line 873 of file contract_core.h.

◆ READ_SPINOR [1/2]

#define READ_SPINOR   READ_SPINOR_DOUBLE_TEX

Definition at line 886 of file contract_core.h.

Referenced by contractGamma5Kernel(), contractKernel(), contractTsliceKernel(), for(), and if().

◆ READ_SPINOR [2/2]

#define READ_SPINOR   READ_SPINOR_SINGLE_TEX

Definition at line 886 of file contract_core.h.

◆ READ_SPINOR_SINGLE

#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)];
sp_idx
int sp_idx
Definition: dw_dslash5inv_core.h:168
spinor
cpuColorSpinorField * spinor
Definition: covdev_test.cpp:41

Definition at line 848 of file contract_core.h.

◆ READ_SPINOR_SINGLE_TEX

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

Definition at line 857 of file contract_core.h.

◆ SPINOR_HOP [1/2]

#define SPINOR_HOP   12

Definition at line 908 of file contract_core.h.

◆ SPINOR_HOP [2/2]

#define SPINOR_HOP   6

Definition at line 908 of file contract_core.h.

◆ SPINORTEX [1/2]

#define SPINORTEX   spinorTexDouble

Definition at line 891 of file contract_core.h.

Referenced by contractGamma5Kernel(), contractKernel(), contractTsliceKernel(), for(), and if().

◆ SPINORTEX [2/2]

#define SPINORTEX   spinorTexSingle

Definition at line 891 of file contract_core.h.

◆ tmp_im

#define tmp_im   tmp.y

Definition at line 9 of file contract_core.h.

Referenced by contractGamma5Kernel(), contractKernel(), and contractTsliceKernel().

◆ tmp_re

#define tmp_re   tmp.x

Definition at line 8 of file contract_core.h.

Referenced by contractGamma5Kernel(), contractKernel(), and contractTsliceKernel().

◆ TOTAL_COMPONENTS

#define TOTAL_COMPONENTS   16

Definition at line 11 of file contract_core.h.

Referenced by contractGamma5Kernel(), contractKernel(), and contractTsliceKernel().

Function Documentation

◆ contractGamma5Kernel() [1/2]

__global__ void contractGamma5Kernel ( double2 *  out,
double2 *  in1,
double2 *  in2,
int  myStride,
const int  Parity,
const DslashParam  param 
)

Definition at line 65 of file contract_core.h.

References blockDim, INTERTEX, out, param, READ_INTERMEDIATE_SPINOR, READ_SPINOR, sid, SPINORTEX, tmp, tmp_im, tmp_re, TOTAL_COMPONENTS, and QudaGaugeParam_s::X.

◆ contractGamma5Kernel() [2/2]

__global__ void contractGamma5Kernel ( float2 *  out,
float4 *  in1,
float4 *  in2,
int  myStride,
const int  Parity,
const DslashParam  param 
)

Definition at line 910 of file contract_core.h.

References blockDim, INTERTEX, out, param, READ_INTERMEDIATE_SPINOR, READ_SPINOR, sid, SPINORTEX, tmp, tmp_im, tmp_re, TOTAL_COMPONENTS, and QudaGaugeParam_s::X.

◆ contractKernel() [1/2]

__global__ void contractKernel ( double2 *  out,
double2 *  in1,
double2 *  in2,
int  myStride,
const int  Parity,
const DslashParam  param 
)

Definition at line 594 of file contract_core.h.

References blockDim, INTERTEX, out, param, READ_INTERMEDIATE_SPINOR, READ_SPINOR, sid, SPINORTEX, tmp, tmp_im, tmp_re, TOTAL_COMPONENTS, and QudaGaugeParam_s::X.

◆ contractKernel() [2/2]

__global__ void contractKernel ( float2 *  out,
float4 *  in1,
float4 *  in2,
int  myStride,
const int  Parity,
const DslashParam  param 
)

Definition at line 1454 of file contract_core.h.

References blockDim, INTERTEX, out, param, READ_INTERMEDIATE_SPINOR, READ_SPINOR, sid, SPINORTEX, tmp, tmp_im, tmp_re, TOTAL_COMPONENTS, and QudaGaugeParam_s::X.

◆ contractTsliceKernel() [1/2]

__global__ void contractTsliceKernel ( double2 *  out,
double2 *  in1,
double2 *  in2,
int  myStride,
const int  Tslice,
const int  Parity,
const DslashParam  param 
)

Definition at line 343 of file contract_core.h.

References blockDim, INTERTEX, out, param, READ_INTERMEDIATE_SPINOR, READ_SPINOR, sid, SPINORTEX, tmp, tmp_im, tmp_re, TOTAL_COMPONENTS, and QudaGaugeParam_s::X.

◆ contractTsliceKernel() [2/2]

__global__ void contractTsliceKernel ( float2 *  out,
float4 *  in1,
float4 *  in2,
int  myStride,
const int  Tslice,
const int  Parity,
const DslashParam  param 
)

Definition at line 1200 of file contract_core.h.

References blockDim, INTERTEX, out, param, READ_INTERMEDIATE_SPINOR, READ_SPINOR, sid, SPINORTEX, tmp, tmp_im, tmp_re, TOTAL_COMPONENTS, and QudaGaugeParam_s::X.

Generated by   doxygen 1.8.14