QUDA  0.9.0
Macros | Functions
contract_core_plus.h File Reference

Go to the source code of this file.

Macros

#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 contractGamma5PlusKernel (double2 *out, double2 *in1, double2 *in2, int myStride, const int Parity, const DslashParam param)
 
__global__ void contractTslicePlusKernel (double2 *out, double2 *in1, double2 *in2, int myStride, const int Tslice, const int Parity, const DslashParam param)
 
__global__ void contractPlusKernel (double2 *out, double2 *in1, double2 *in2, int myStride, const int Parity, const DslashParam param)
 
__global__ void contractGamma5PlusKernel (float2 *out, float4 *in1, float4 *in2, int myStride, const int Parity, const DslashParam param)
 
__global__ void contractTslicePlusKernel (float2 *out, float4 *in1, float4 *in2, int myStride, const int Tslice, const int Parity, const DslashParam param)
 
__global__ void contractPlusKernel (float2 *out, float4 *in1, float4 *in2, int myStride, const int Parity, const DslashParam param)
 

Macro Definition Documentation

◆ INTERTEX [1/2]

#define INTERTEX   interTexDouble

Definition at line 947 of file contract_core_plus.h.

Referenced by contractGamma5PlusKernel(), contractPlusKernel(), and contractTslicePlusKernel().

◆ INTERTEX [2/2]

#define INTERTEX   interTexSingle

Definition at line 947 of file contract_core_plus.h.

◆ READ_INTERMEDIATE_SPINOR [1/2]

#define READ_INTERMEDIATE_SPINOR   READ_INTERMEDIATE_SPINOR_DOUBLE_TEX

Definition at line 942 of file contract_core_plus.h.

Referenced by contractGamma5PlusKernel(), contractPlusKernel(), and contractTslicePlusKernel().

◆ READ_INTERMEDIATE_SPINOR [2/2]

#define READ_INTERMEDIATE_SPINOR   READ_INTERMEDIATE_SPINOR_SINGLE_TEX

Definition at line 942 of file contract_core_plus.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 9 of file contract_core_plus.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 23 of file contract_core_plus.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 908 of file contract_core_plus.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 916 of file contract_core_plus.h.

◆ READ_SPINOR [1/2]

#define READ_SPINOR   READ_SPINOR_DOUBLE_TEX

Definition at line 929 of file contract_core_plus.h.

Referenced by contractGamma5PlusKernel(), contractPlusKernel(), and contractTslicePlusKernel().

◆ READ_SPINOR [2/2]

#define READ_SPINOR   READ_SPINOR_SINGLE_TEX

Definition at line 929 of file contract_core_plus.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 891 of file contract_core_plus.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 900 of file contract_core_plus.h.

◆ SPINOR_HOP [1/2]

#define SPINOR_HOP   12

Definition at line 951 of file contract_core_plus.h.

◆ SPINOR_HOP [2/2]

#define SPINOR_HOP   6

Definition at line 951 of file contract_core_plus.h.

◆ SPINORTEX [1/2]

#define SPINORTEX   spinorTexDouble

Definition at line 934 of file contract_core_plus.h.

Referenced by contractGamma5PlusKernel(), contractPlusKernel(), and contractTslicePlusKernel().

◆ SPINORTEX [2/2]

#define SPINORTEX   spinorTexSingle

Definition at line 934 of file contract_core_plus.h.

◆ tmp_im

#define tmp_im   tmp.y

Definition at line 5 of file contract_core_plus.h.

Referenced by contractGamma5PlusKernel(), contractPlusKernel(), and contractTslicePlusKernel().

◆ tmp_re

#define tmp_re   tmp.x

Definition at line 4 of file contract_core_plus.h.

Referenced by contractGamma5PlusKernel(), contractPlusKernel(), and contractTslicePlusKernel().

◆ TOTAL_COMPONENTS

#define TOTAL_COMPONENTS   16

Definition at line 7 of file contract_core_plus.h.

Referenced by contractGamma5PlusKernel(), contractPlusKernel(), and contractTslicePlusKernel().

Function Documentation

◆ contractGamma5PlusKernel() [1/2]

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

Definition at line 58 of file contract_core_plus.h.

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

◆ contractGamma5PlusKernel() [2/2]

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

Definition at line 953 of file contract_core_plus.h.

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

◆ contractPlusKernel() [1/2]

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

Definition at line 621 of file contract_core_plus.h.

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

◆ contractPlusKernel() [2/2]

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

Definition at line 1531 of file contract_core_plus.h.

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

◆ contractTslicePlusKernel() [1/2]

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

Definition at line 353 of file contract_core_plus.h.

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

◆ contractTslicePlusKernel() [2/2]

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

Definition at line 1260 of file contract_core_plus.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