QUDA  0.9.0
Macros | Functions | Variables
dw_dslash_core.h File Reference
#include "read_gauge.h"
#include "io_spinor.h"
Include dependency graph for dw_dslash_core.h:

Go to the source code of this file.

Macros

#define DSLASH_SHARED_FLOATS_PER_THREAD   0
 
#define VOLATILE   volatile
 
#define spinorFloat   float
 
#define i00_re   I0.x
 
#define i00_im   I0.y
 
#define i01_re   I0.z
 
#define i01_im   I0.w
 
#define i02_re   I1.x
 
#define i02_im   I1.y
 
#define i10_re   I1.z
 
#define i10_im   I1.w
 
#define i11_re   I2.x
 
#define i11_im   I2.y
 
#define i12_re   I2.z
 
#define i12_im   I2.w
 
#define i20_re   I3.x
 
#define i20_im   I3.y
 
#define i21_re   I3.z
 
#define i21_im   I3.w
 
#define i22_re   I4.x
 
#define i22_im   I4.y
 
#define i30_re   I4.z
 
#define i30_im   I4.w
 
#define i31_re   I5.x
 
#define i31_im   I5.y
 
#define i32_re   I5.z
 
#define i32_im   I5.w
 
#define g00_re   G0.x
 
#define g00_im   G0.y
 
#define g01_re   G0.z
 
#define g01_im   G0.w
 
#define g02_re   G1.x
 
#define g02_im   G1.y
 
#define g10_re   G1.z
 
#define g10_im   G1.w
 
#define g11_re   G2.x
 
#define g11_im   G2.y
 
#define g12_re   G2.z
 
#define g12_im   G2.w
 
#define g20_re   G3.x
 
#define g20_im   G3.y
 
#define g21_re   G3.z
 
#define g21_im   G3.w
 
#define g22_re   G4.x
 
#define g22_im   G4.y
 
#define gT00_re   (+g00_re)
 
#define gT00_im   (-g00_im)
 
#define gT01_re   (+g10_re)
 
#define gT01_im   (-g10_im)
 
#define gT02_re   (+g20_re)
 
#define gT02_im   (-g20_im)
 
#define gT10_re   (+g01_re)
 
#define gT10_im   (-g01_im)
 
#define gT11_re   (+g11_re)
 
#define gT11_im   (-g11_im)
 
#define gT12_re   (+g21_re)
 
#define gT12_im   (-g21_im)
 
#define gT20_re   (+g02_re)
 
#define gT20_im   (-g02_im)
 
#define gT21_re   (+g12_re)
 
#define gT21_im   (-g12_im)
 
#define gT22_re   (+g22_re)
 
#define gT22_im   (-g22_im)
 
#define SHARED_STRIDE   16
 

Functions

 if (sid >=param.threads *param.dc.Ls) return
 
 coordsFromIndex< 5, QUDA_5D_PC, EVEN_X > (X, coord, sid, param)
 
 if (! s_parity)
 
 READ_SPINOR (SPINORTEX, param.sp_stride, sp_idx, sp_idx)
 
 RECONSTRUCT_GAUGE_MATRIX (0)
 
 RECONSTRUCT_GAUGE_MATRIX (1)
 
 RECONSTRUCT_GAUGE_MATRIX (2)
 
 RECONSTRUCT_GAUGE_MATRIX (3)
 
 RECONSTRUCT_GAUGE_MATRIX (4)
 
 RECONSTRUCT_GAUGE_MATRIX (5)
 
 if (param.gauge_fixed &&ga_idx< param.dc.X4X3X2X1hmX3X2X1h)
 
 READ_SPINOR_DOWN (SPINORTEX, param.sp_stride, sp_idx, sp_idx)
 
 RECONSTRUCT_GAUGE_MATRIX (6)
 
 READ_SPINOR_UP (SPINORTEX, param.sp_stride, sp_idx, sp_idx)
 
 RECONSTRUCT_GAUGE_MATRIX (7)
 

Variables

VOLATILE spinorFloat o00_re = 0
 
VOLATILE spinorFloat o00_im = 0
 
VOLATILE spinorFloat o01_re = 0
 
VOLATILE spinorFloat o01_im = 0
 
VOLATILE spinorFloat o02_re = 0
 
VOLATILE spinorFloat o02_im = 0
 
VOLATILE spinorFloat o10_re = 0
 
VOLATILE spinorFloat o10_im = 0
 
VOLATILE spinorFloat o11_re = 0
 
VOLATILE spinorFloat o11_im = 0
 
VOLATILE spinorFloat o12_re = 0
 
VOLATILE spinorFloat o12_im = 0
 
VOLATILE spinorFloat o20_re = 0
 
VOLATILE spinorFloat o20_im = 0
 
VOLATILE spinorFloat o21_re = 0
 
VOLATILE spinorFloat o21_im = 0
 
VOLATILE spinorFloat o22_re = 0
 
VOLATILE spinorFloat o22_im = 0
 
VOLATILE spinorFloat o30_re = 0
 
VOLATILE spinorFloat o30_im = 0
 
VOLATILE spinorFloat o31_re = 0
 
VOLATILE spinorFloat o31_im = 0
 
VOLATILE spinorFloat o32_re = 0
 
VOLATILE spinorFloat o32_im = 0
 
int sid = ((blockIdx.y*blockDim.y + threadIdx.y)*gridDim.x + blockIdx.x)*blockDim.x + threadIdx.x
 
int X
 
int coord [5]
 
int s_parity = ( sid/param.dc.volume_4d_cb ) % 2
 
float4 G0
 
float4 G1
 
float4 G2
 
float4 G3
 
float4 G4
 
const int ga_idx = sid % param.dc.volume_4d_cb
 
 else { ASSN_GAUGE_MATRIX(G, GAUGE1TEX, 0, ga_idx, param.gauge_stride)
 
spinorFloat a0_re = +i00_re+i30_im
 
spinorFloat a0_im = +i00_im-i30_re
 
spinorFloat a1_re = +i01_re+i31_im
 
spinorFloat a1_im = +i01_im-i31_re
 
spinorFloat a2_re = +i02_re+i32_im
 
spinorFloat a2_im = +i02_im-i32_re
 
spinorFloat b0_re = +i10_re+i20_im
 
spinorFloat b0_im = +i10_im-i20_re
 
spinorFloat b1_re = +i11_re+i21_im
 
spinorFloat b1_im = +i11_im-i21_re
 
spinorFloat b2_re = +i12_re+i22_im
 
spinorFloat b2_im = +i12_im-i22_re
 
spinorFloat A0_re = 0
 
spinorFloat A0_im = 0
 
spinorFloat B0_re = 0
 
spinorFloat B0_im = 0
 
spinorFloat A1_re = 0
 
spinorFloat A1_im = 0
 
spinorFloat B1_re = 0
 
spinorFloat B1_im = 0
 
spinorFloat A2_re = 0
 
spinorFloat A2_im = 0
 
spinorFloat B2_re = 0
 
spinorFloat B2_im = 0
 

Macro Definition Documentation

◆ DSLASH_SHARED_FLOATS_PER_THREAD

#define DSLASH_SHARED_FLOATS_PER_THREAD   0

Definition at line 3 of file dw_dslash_core.h.

◆ g00_im

#define g00_im   G0.y

Definition at line 90 of file dw_dslash_core.h.

◆ g00_re

#define g00_re   G0.x

Definition at line 89 of file dw_dslash_core.h.

◆ g01_im

#define g01_im   G0.w

Definition at line 92 of file dw_dslash_core.h.

◆ g01_re

#define g01_re   G0.z

Definition at line 91 of file dw_dslash_core.h.

◆ g02_im

#define g02_im   G1.y

Definition at line 94 of file dw_dslash_core.h.

◆ g02_re

#define g02_re   G1.x

Definition at line 93 of file dw_dslash_core.h.

◆ g10_im

#define g10_im   G1.w

Definition at line 96 of file dw_dslash_core.h.

◆ g10_re

#define g10_re   G1.z

Definition at line 95 of file dw_dslash_core.h.

◆ g11_im

#define g11_im   G2.y

Definition at line 98 of file dw_dslash_core.h.

◆ g11_re

#define g11_re   G2.x

Definition at line 97 of file dw_dslash_core.h.

◆ g12_im

#define g12_im   G2.w

Definition at line 100 of file dw_dslash_core.h.

◆ g12_re

#define g12_re   G2.z

Definition at line 99 of file dw_dslash_core.h.

◆ g20_im

#define g20_im   G3.y

Definition at line 102 of file dw_dslash_core.h.

◆ g20_re

#define g20_re   G3.x

Definition at line 101 of file dw_dslash_core.h.

◆ g21_im

#define g21_im   G3.w

Definition at line 104 of file dw_dslash_core.h.

◆ g21_re

#define g21_re   G3.z

Definition at line 103 of file dw_dslash_core.h.

◆ g22_im

#define g22_im   G4.y

Definition at line 106 of file dw_dslash_core.h.

◆ g22_re

#define g22_re   G4.x

Definition at line 105 of file dw_dslash_core.h.

◆ gT00_im

#define gT00_im   (-g00_im)

Definition at line 112 of file dw_dslash_core.h.

◆ gT00_re

#define gT00_re   (+g00_re)

Definition at line 111 of file dw_dslash_core.h.

◆ gT01_im

#define gT01_im   (-g10_im)

Definition at line 114 of file dw_dslash_core.h.

◆ gT01_re

#define gT01_re   (+g10_re)

Definition at line 113 of file dw_dslash_core.h.

◆ gT02_im

#define gT02_im   (-g20_im)

Definition at line 116 of file dw_dslash_core.h.

◆ gT02_re

#define gT02_re   (+g20_re)

Definition at line 115 of file dw_dslash_core.h.

◆ gT10_im

#define gT10_im   (-g01_im)

Definition at line 118 of file dw_dslash_core.h.

◆ gT10_re

#define gT10_re   (+g01_re)

Definition at line 117 of file dw_dslash_core.h.

◆ gT11_im

#define gT11_im   (-g11_im)

Definition at line 120 of file dw_dslash_core.h.

◆ gT11_re

#define gT11_re   (+g11_re)

Definition at line 119 of file dw_dslash_core.h.

◆ gT12_im

#define gT12_im   (-g21_im)

Definition at line 122 of file dw_dslash_core.h.

◆ gT12_re

#define gT12_re   (+g21_re)

Definition at line 121 of file dw_dslash_core.h.

◆ gT20_im

#define gT20_im   (-g02_im)

Definition at line 124 of file dw_dslash_core.h.

◆ gT20_re

#define gT20_re   (+g02_re)

Definition at line 123 of file dw_dslash_core.h.

◆ gT21_im

#define gT21_im   (-g12_im)

Definition at line 126 of file dw_dslash_core.h.

◆ gT21_re

#define gT21_re   (+g12_re)

Definition at line 125 of file dw_dslash_core.h.

◆ gT22_im

#define gT22_im   (-g22_im)

Definition at line 128 of file dw_dslash_core.h.

◆ gT22_re

#define gT22_re   (+g22_re)

Definition at line 127 of file dw_dslash_core.h.

◆ i00_im

#define i00_im   I0.y

Definition at line 42 of file dw_dslash_core.h.

Referenced by if().

◆ i00_re

#define i00_re   I0.x

Definition at line 41 of file dw_dslash_core.h.

Referenced by if().

◆ i01_im

#define i01_im   I0.w

Definition at line 44 of file dw_dslash_core.h.

Referenced by if().

◆ i01_re

#define i01_re   I0.z

Definition at line 43 of file dw_dslash_core.h.

Referenced by if().

◆ i02_im

#define i02_im   I1.y

Definition at line 46 of file dw_dslash_core.h.

Referenced by if().

◆ i02_re

#define i02_re   I1.x

Definition at line 45 of file dw_dslash_core.h.

Referenced by if().

◆ i10_im

#define i10_im   I1.w

Definition at line 48 of file dw_dslash_core.h.

Referenced by if().

◆ i10_re

#define i10_re   I1.z

Definition at line 47 of file dw_dslash_core.h.

Referenced by if().

◆ i11_im

#define i11_im   I2.y

Definition at line 50 of file dw_dslash_core.h.

Referenced by if().

◆ i11_re

#define i11_re   I2.x

Definition at line 49 of file dw_dslash_core.h.

Referenced by if().

◆ i12_im

#define i12_im   I2.w

Definition at line 52 of file dw_dslash_core.h.

Referenced by if().

◆ i12_re

#define i12_re   I2.z

Definition at line 51 of file dw_dslash_core.h.

Referenced by if().

◆ i20_im

#define i20_im   I3.y

Definition at line 54 of file dw_dslash_core.h.

Referenced by if().

◆ i20_re

#define i20_re   I3.x

Definition at line 53 of file dw_dslash_core.h.

Referenced by if().

◆ i21_im

#define i21_im   I3.w

Definition at line 56 of file dw_dslash_core.h.

Referenced by if().

◆ i21_re

#define i21_re   I3.z

Definition at line 55 of file dw_dslash_core.h.

Referenced by if().

◆ i22_im

#define i22_im   I4.y

Definition at line 58 of file dw_dslash_core.h.

Referenced by if().

◆ i22_re

#define i22_re   I4.x

Definition at line 57 of file dw_dslash_core.h.

Referenced by if().

◆ i30_im

#define i30_im   I4.w

Definition at line 60 of file dw_dslash_core.h.

Referenced by if().

◆ i30_re

#define i30_re   I4.z

Definition at line 59 of file dw_dslash_core.h.

Referenced by if().

◆ i31_im

#define i31_im   I5.y

Definition at line 62 of file dw_dslash_core.h.

Referenced by if().

◆ i31_re

#define i31_re   I5.x

Definition at line 61 of file dw_dslash_core.h.

Referenced by if().

◆ i32_im

#define i32_im   I5.w

Definition at line 64 of file dw_dslash_core.h.

Referenced by if().

◆ i32_re

#define i32_re   I5.z

Definition at line 63 of file dw_dslash_core.h.

Referenced by if().

◆ SHARED_STRIDE

#define SHARED_STRIDE   16

Definition at line 166 of file dw_dslash_core.h.

◆ spinorFloat

#define spinorFloat   float

Definition at line 40 of file dw_dslash_core.h.

Referenced by if().

◆ VOLATILE

#define VOLATILE   volatile

Definition at line 10 of file dw_dslash_core.h.

Function Documentation

◆ coordsFromIndex< 5, QUDA_5D_PC, EVEN_X >()

coordsFromIndex< 5, QUDA_5D_PC, EVEN_X > ( X  ,
coord  ,
sid  ,
param   
)

◆ if() [1/3]

if ( sid >=param.threads *param.dc.  Ls)

◆ if() [2/3]

if ( s_parity)

Definition at line 291 of file dw_dslash_core.h.

◆ if() [3/3]

if ( )

◆ READ_SPINOR()

READ_SPINOR ( SPINORTEX  ,
param.  sp_stride,
sp_idx  ,
sp_idx   
)

◆ READ_SPINOR_DOWN()

READ_SPINOR_DOWN ( SPINORTEX  ,
param.  sp_stride,
sp_idx  ,
sp_idx   
)

Referenced by if().

Here is the caller graph for this function:

◆ READ_SPINOR_UP()

READ_SPINOR_UP ( SPINORTEX  ,
param.  sp_stride,
sp_idx  ,
sp_idx   
)

◆ RECONSTRUCT_GAUGE_MATRIX() [1/8]

RECONSTRUCT_GAUGE_MATRIX ( )

◆ RECONSTRUCT_GAUGE_MATRIX() [2/8]

RECONSTRUCT_GAUGE_MATRIX ( )

◆ RECONSTRUCT_GAUGE_MATRIX() [3/8]

RECONSTRUCT_GAUGE_MATRIX ( )

◆ RECONSTRUCT_GAUGE_MATRIX() [4/8]

RECONSTRUCT_GAUGE_MATRIX ( )

◆ RECONSTRUCT_GAUGE_MATRIX() [5/8]

RECONSTRUCT_GAUGE_MATRIX ( )

◆ RECONSTRUCT_GAUGE_MATRIX() [6/8]

RECONSTRUCT_GAUGE_MATRIX ( )

◆ RECONSTRUCT_GAUGE_MATRIX() [7/8]

RECONSTRUCT_GAUGE_MATRIX ( )

◆ RECONSTRUCT_GAUGE_MATRIX() [8/8]

RECONSTRUCT_GAUGE_MATRIX ( )

Variable Documentation

◆ a0_im

a0_im = +i00_im-i30_re

Definition at line 294 of file dw_dslash_core.h.

Referenced by if().

◆ A0_im

A0_im = 0

Definition at line 351 of file dw_dslash_core.h.

Referenced by if().

◆ a0_re

a0_re = +i00_re+i30_im

Definition at line 294 of file dw_dslash_core.h.

Referenced by if().

◆ A0_re

A0_re = 0

Definition at line 344 of file dw_dslash_core.h.

Referenced by if().

◆ a1_im

a1_im = +i01_im-i31_re

Definition at line 295 of file dw_dslash_core.h.

Referenced by if().

◆ A1_im

A1_im = 0

Definition at line 381 of file dw_dslash_core.h.

Referenced by if().

◆ a1_re

a1_re = +i01_re+i31_im

Definition at line 295 of file dw_dslash_core.h.

Referenced by if().

◆ A1_re

A1_re = 0

Definition at line 374 of file dw_dslash_core.h.

Referenced by if().

◆ a2_im

a2_im = +i02_im-i32_re

Definition at line 296 of file dw_dslash_core.h.

Referenced by if().

◆ A2_im

A2_im = 0

Definition at line 411 of file dw_dslash_core.h.

Referenced by if().

◆ a2_re

a2_re = +i02_re+i32_im

Definition at line 296 of file dw_dslash_core.h.

Referenced by if().

◆ A2_re

A2_re = 0

Definition at line 404 of file dw_dslash_core.h.

Referenced by if().

◆ b0_im

b0_im = +i10_im-i20_re

Definition at line 297 of file dw_dslash_core.h.

Referenced by if().

◆ B0_im

B0_im = 0

Definition at line 365 of file dw_dslash_core.h.

Referenced by if().

◆ b0_re

b0_re = +i10_re+i20_im

Definition at line 297 of file dw_dslash_core.h.

Referenced by if().

◆ B0_re

B0_re = 0

Definition at line 358 of file dw_dslash_core.h.

Referenced by if().

◆ b1_im

b1_im = +i11_im-i21_re

Definition at line 298 of file dw_dslash_core.h.

Referenced by if().

◆ B1_im

B1_im = 0

Definition at line 395 of file dw_dslash_core.h.

Referenced by if().

◆ b1_re

b1_re = +i11_re+i21_im

Definition at line 298 of file dw_dslash_core.h.

Referenced by if().

◆ B1_re

B1_re = 0

Definition at line 388 of file dw_dslash_core.h.

Referenced by if().

◆ b2_im

b2_im = +i12_im-i22_re

Definition at line 299 of file dw_dslash_core.h.

Referenced by if().

◆ B2_im

B2_im = 0

Definition at line 425 of file dw_dslash_core.h.

Referenced by if().

◆ b2_re

b2_re = +i12_re+i22_im

Definition at line 299 of file dw_dslash_core.h.

Referenced by if().

◆ B2_re

B2_re = 0

Definition at line 418 of file dw_dslash_core.h.

Referenced by if().

◆ coord

int coord[5]

Definition at line 177 of file dw_dslash_core.h.

◆ else

else { ASSN_GAUGE_MATRIX(G, GAUGE1TEX, 0, ga_idx, param.gauge_stride)

Definition at line 292 of file dw_dslash_core.h.

◆ G0

float4 G0

Definition at line 258 of file dw_dslash_core.h.

◆ G1

float4 G1

Definition at line 259 of file dw_dslash_core.h.

◆ G2

float4 G2

Definition at line 260 of file dw_dslash_core.h.

◆ G3

float4 G3

Definition at line 261 of file dw_dslash_core.h.

◆ G4

float4 G4

Definition at line 262 of file dw_dslash_core.h.

◆ ga_idx

const int ga_idx = sid % param.dc.volume_4d_cb

Definition at line 288 of file dw_dslash_core.h.

Referenced by if().

◆ o00_im

o00_im = 0

Definition at line 132 of file dw_dslash_core.h.

◆ o00_re

o00_re = 0

Definition at line 131 of file dw_dslash_core.h.

◆ o01_im

o01_im = 0

Definition at line 134 of file dw_dslash_core.h.

◆ o01_re

o01_re = 0

Definition at line 133 of file dw_dslash_core.h.

◆ o02_im

o02_im = 0

Definition at line 136 of file dw_dslash_core.h.

◆ o02_re

o02_re = 0

Definition at line 135 of file dw_dslash_core.h.

◆ o10_im

o10_im = 0

Definition at line 138 of file dw_dslash_core.h.

◆ o10_re

o10_re = 0

Definition at line 137 of file dw_dslash_core.h.

◆ o11_im

o11_im = 0

Definition at line 140 of file dw_dslash_core.h.

◆ o11_re

o11_re = 0

Definition at line 139 of file dw_dslash_core.h.

◆ o12_im

o12_im = 0

Definition at line 142 of file dw_dslash_core.h.

◆ o12_re

o12_re = 0

Definition at line 141 of file dw_dslash_core.h.

◆ o20_im

o20_im = 0

Definition at line 144 of file dw_dslash_core.h.

Referenced by if().

◆ o20_re

o20_re = 0

Definition at line 143 of file dw_dslash_core.h.

Referenced by if().

◆ o21_im

o21_im = 0

Definition at line 146 of file dw_dslash_core.h.

Referenced by if().

◆ o21_re

o21_re = 0

Definition at line 145 of file dw_dslash_core.h.

Referenced by if().

◆ o22_im

o22_im = 0

Definition at line 148 of file dw_dslash_core.h.

Referenced by if().

◆ o22_re

o22_re = 0

Definition at line 147 of file dw_dslash_core.h.

Referenced by if().

◆ o30_im

o30_im = 0

Definition at line 150 of file dw_dslash_core.h.

Referenced by if().

◆ o30_re

o30_re = 0

Definition at line 149 of file dw_dslash_core.h.

Referenced by if().

◆ o31_im

o31_im = 0

Definition at line 152 of file dw_dslash_core.h.

Referenced by if().

◆ o31_re

o31_re = 0

Definition at line 151 of file dw_dslash_core.h.

Referenced by if().

◆ o32_im

o32_im = 0

Definition at line 154 of file dw_dslash_core.h.

Referenced by if().

◆ o32_re

o32_re = 0

Definition at line 153 of file dw_dslash_core.h.

Referenced by if().

◆ s_parity

s_parity = ( sid/param.dc.volume_4d_cb ) % 2

Definition at line 179 of file dw_dslash_core.h.

◆ sid

Definition at line 174 of file dw_dslash_core.h.

◆ X

int X

Definition at line 177 of file dw_dslash_core.h.