dw_dslash_dagger_core.h File Reference

#include "read_gauge.h"
#include "io_spinor.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.Ls) return
	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 (gauge_fixed &&ga_idx< X4X3X2X1hmX3X2X1h)
	READ_SPINOR_UP (SPINORTEX, param.sp_stride, sp_idx, sp_idx)
	RECONSTRUCT_GAUGE_MATRIX (6)
	READ_SPINOR_DOWN (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 = 2*sid + (boundaryCrossing + param.parity) % 2
int	x1 = X % X1
int	x2 = (X/X1) % X2
int	x3 = (X/(X1*X2)) % X3
int	x4 = (X/(X1*X2*X3)) % X4
int	xs = X/(X1*X2*X3*X4)
int	s_parity = ( sid/(X4*X3*X2*X1h) ) % 2
int	boundaryCrossing = sid/X1h + sid/(X2*X1h) + sid/(X3*X2*X1h) + sid/(X4*X3*X2*X1h)
float4	G0
float4	G1
float4	G2
float4	G3
float4	G4
const int	ga_idx = sid % Vh
	else { ASSN_GAUGE_MATRIX(G, GAUGE1TEX, 0, ga_idx, ga_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

#define DSLASH_SHARED_FLOATS_PER_THREAD   0

Definition at line 3 of file dw_dslash_dagger_core.h.

#define g00_im   G0.y

Definition at line 90 of file dw_dslash_dagger_core.h.

#define g00_re   G0.x

Definition at line 89 of file dw_dslash_dagger_core.h.

#define g01_im   G0.w

Definition at line 92 of file dw_dslash_dagger_core.h.

#define g01_re   G0.z

Definition at line 91 of file dw_dslash_dagger_core.h.

#define g02_im   G1.y

Definition at line 94 of file dw_dslash_dagger_core.h.

#define g02_re   G1.x

Definition at line 93 of file dw_dslash_dagger_core.h.

#define g10_im   G1.w

Definition at line 96 of file dw_dslash_dagger_core.h.

#define g10_re   G1.z

Definition at line 95 of file dw_dslash_dagger_core.h.

#define g11_im   G2.y

Definition at line 98 of file dw_dslash_dagger_core.h.

#define g11_re   G2.x

Definition at line 97 of file dw_dslash_dagger_core.h.

#define g12_im   G2.w

Definition at line 100 of file dw_dslash_dagger_core.h.

#define g12_re   G2.z

Definition at line 99 of file dw_dslash_dagger_core.h.

#define g20_im   G3.y

Definition at line 102 of file dw_dslash_dagger_core.h.

#define g20_re   G3.x

Definition at line 101 of file dw_dslash_dagger_core.h.

#define g21_im   G3.w

Definition at line 104 of file dw_dslash_dagger_core.h.

#define g21_re   G3.z

Definition at line 103 of file dw_dslash_dagger_core.h.

#define g22_im   G4.y

Definition at line 106 of file dw_dslash_dagger_core.h.

#define g22_re   G4.x

Definition at line 105 of file dw_dslash_dagger_core.h.

#define gT00_im   (-g00_im)

Definition at line 112 of file dw_dslash_dagger_core.h.

#define gT00_re   (+g00_re)

Definition at line 111 of file dw_dslash_dagger_core.h.

#define gT01_im   (-g10_im)

Definition at line 114 of file dw_dslash_dagger_core.h.

#define gT01_re   (+g10_re)

Definition at line 113 of file dw_dslash_dagger_core.h.

#define gT02_im   (-g20_im)

Definition at line 116 of file dw_dslash_dagger_core.h.

#define gT02_re   (+g20_re)

Definition at line 115 of file dw_dslash_dagger_core.h.

#define gT10_im   (-g01_im)

Definition at line 118 of file dw_dslash_dagger_core.h.

#define gT10_re   (+g01_re)

Definition at line 117 of file dw_dslash_dagger_core.h.

#define gT11_im   (-g11_im)

Definition at line 120 of file dw_dslash_dagger_core.h.

#define gT11_re   (+g11_re)

Definition at line 119 of file dw_dslash_dagger_core.h.

#define gT12_im   (-g21_im)

Definition at line 122 of file dw_dslash_dagger_core.h.

#define gT12_re   (+g21_re)

Definition at line 121 of file dw_dslash_dagger_core.h.

#define gT20_im   (-g02_im)

Definition at line 124 of file dw_dslash_dagger_core.h.

#define gT20_re   (+g02_re)

Definition at line 123 of file dw_dslash_dagger_core.h.

#define gT21_im   (-g12_im)

Definition at line 126 of file dw_dslash_dagger_core.h.

#define gT21_re   (+g12_re)

Definition at line 125 of file dw_dslash_dagger_core.h.

#define gT22_im   (-g22_im)

Definition at line 128 of file dw_dslash_dagger_core.h.

#define gT22_re   (+g22_re)

Definition at line 127 of file dw_dslash_dagger_core.h.

#define i00_im   I0.y

Definition at line 42 of file dw_dslash_dagger_core.h.

#define i00_re   I0.x

Definition at line 41 of file dw_dslash_dagger_core.h.

#define i01_im   I0.w

Definition at line 44 of file dw_dslash_dagger_core.h.

#define i01_re   I0.z

Definition at line 43 of file dw_dslash_dagger_core.h.

#define i02_im   I1.y

Definition at line 46 of file dw_dslash_dagger_core.h.

#define i02_re   I1.x

Definition at line 45 of file dw_dslash_dagger_core.h.

#define i10_im   I1.w

Definition at line 48 of file dw_dslash_dagger_core.h.

#define i10_re   I1.z

Definition at line 47 of file dw_dslash_dagger_core.h.

#define i11_im   I2.y

Definition at line 50 of file dw_dslash_dagger_core.h.

#define i11_re   I2.x

Definition at line 49 of file dw_dslash_dagger_core.h.

#define i12_im   I2.w

Definition at line 52 of file dw_dslash_dagger_core.h.

#define i12_re   I2.z

Definition at line 51 of file dw_dslash_dagger_core.h.

#define i20_im   I3.y

Definition at line 54 of file dw_dslash_dagger_core.h.

#define i20_re   I3.x

Definition at line 53 of file dw_dslash_dagger_core.h.

#define i21_im   I3.w

Definition at line 56 of file dw_dslash_dagger_core.h.

#define i21_re   I3.z

Definition at line 55 of file dw_dslash_dagger_core.h.

#define i22_im   I4.y

Definition at line 58 of file dw_dslash_dagger_core.h.

#define i22_re   I4.x

Definition at line 57 of file dw_dslash_dagger_core.h.

#define i30_im   I4.w

Definition at line 60 of file dw_dslash_dagger_core.h.

#define i30_re   I4.z

Definition at line 59 of file dw_dslash_dagger_core.h.

#define i31_im   I5.y

Definition at line 62 of file dw_dslash_dagger_core.h.

#define i31_re   I5.x

Definition at line 61 of file dw_dslash_dagger_core.h.

#define i32_im   I5.w

Definition at line 64 of file dw_dslash_dagger_core.h.

#define i32_re   I5.z

Definition at line 63 of file dw_dslash_dagger_core.h.

#define SHARED_STRIDE   16

Definition at line 166 of file dw_dslash_dagger_core.h.

#define spinorFloat   float

Definition at line 40 of file dw_dslash_dagger_core.h.

#define VOLATILE   volatile

Definition at line 10 of file dw_dslash_dagger_core.h.

Function Documentation

if

(

sid >=param.threads *param.

Ls

)

if

(

!

s_parity

)

Definition at line 310 of file dw_dslash_dagger_core.h.

if

(

)

Definition at line 1467 of file dw_dslash_dagger_core.h.

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

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

READ_SPINOR_UP	(	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

)

RECONSTRUCT_GAUGE_MATRIX

(

6

)

RECONSTRUCT_GAUGE_MATRIX

(

7

)

Variable Documentation

a0_im = +i00_im+i30_re

Definition at line 313 of file dw_dslash_dagger_core.h.

A0_im = 0

Definition at line 370 of file dw_dslash_dagger_core.h.

a0_re = +i00_re-i30_im

Definition at line 313 of file dw_dslash_dagger_core.h.

A0_re = 0

Definition at line 363 of file dw_dslash_dagger_core.h.

a1_im = +i01_im+i31_re

Definition at line 314 of file dw_dslash_dagger_core.h.

A1_im = 0

Definition at line 400 of file dw_dslash_dagger_core.h.

a1_re = +i01_re-i31_im

Definition at line 314 of file dw_dslash_dagger_core.h.

A1_re = 0

Definition at line 393 of file dw_dslash_dagger_core.h.

a2_im = +i02_im+i32_re

Definition at line 315 of file dw_dslash_dagger_core.h.

A2_im = 0

Definition at line 430 of file dw_dslash_dagger_core.h.

a2_re = +i02_re-i32_im

Definition at line 315 of file dw_dslash_dagger_core.h.

A2_re = 0

Definition at line 423 of file dw_dslash_dagger_core.h.

b0_im = +i10_im+i20_re

Definition at line 316 of file dw_dslash_dagger_core.h.

B0_im = 0

Definition at line 384 of file dw_dslash_dagger_core.h.

b0_re = +i10_re-i20_im

Definition at line 316 of file dw_dslash_dagger_core.h.

B0_re = 0

Definition at line 377 of file dw_dslash_dagger_core.h.

b1_im = +i11_im+i21_re

Definition at line 317 of file dw_dslash_dagger_core.h.

B1_im = 0

Definition at line 414 of file dw_dslash_dagger_core.h.

b1_re = +i11_re-i21_im

Definition at line 317 of file dw_dslash_dagger_core.h.

B1_re = 0

Definition at line 407 of file dw_dslash_dagger_core.h.

b2_im = +i12_im+i22_re

Definition at line 318 of file dw_dslash_dagger_core.h.

B2_im = 0

Definition at line 444 of file dw_dslash_dagger_core.h.

b2_re = +i12_re-i22_im

Definition at line 318 of file dw_dslash_dagger_core.h.

B2_re = 0

Definition at line 437 of file dw_dslash_dagger_core.h.

boundaryCrossing = sid/X1h + sid/(X2*X1h) + sid/(X3*X2*X1h) + sid/(X4*X3*X2*X1h)

Definition at line 182 of file dw_dslash_dagger_core.h.

else { ASSN_GAUGE_MATRIX(G, GAUGE1TEX, 0, ga_idx, ga_stride)

Definition at line 311 of file dw_dslash_dagger_core.h.

float4 G0

Definition at line 280 of file dw_dslash_dagger_core.h.

float4 G1

Definition at line 281 of file dw_dslash_dagger_core.h.

float4 G2

Definition at line 282 of file dw_dslash_dagger_core.h.

float4 G3

Definition at line 283 of file dw_dslash_dagger_core.h.

float4 G4

Definition at line 284 of file dw_dslash_dagger_core.h.

const int ga_idx = sid % Vh

Definition at line 307 of file dw_dslash_dagger_core.h.

o00_im = 0

Definition at line 132 of file dw_dslash_dagger_core.h.

o00_re = 0

Definition at line 131 of file dw_dslash_dagger_core.h.

o01_im = 0

Definition at line 134 of file dw_dslash_dagger_core.h.

o01_re = 0

Definition at line 133 of file dw_dslash_dagger_core.h.

o02_im = 0

Definition at line 136 of file dw_dslash_dagger_core.h.

o02_re = 0

Definition at line 135 of file dw_dslash_dagger_core.h.

o10_im = 0

Definition at line 138 of file dw_dslash_dagger_core.h.

o10_re = 0

Definition at line 137 of file dw_dslash_dagger_core.h.

o11_im = 0

Definition at line 140 of file dw_dslash_dagger_core.h.

o11_re = 0

Definition at line 139 of file dw_dslash_dagger_core.h.

o12_im = 0

Definition at line 142 of file dw_dslash_dagger_core.h.

o12_re = 0

Definition at line 141 of file dw_dslash_dagger_core.h.

o20_im = 0

Definition at line 144 of file dw_dslash_dagger_core.h.

o20_re = 0

Definition at line 143 of file dw_dslash_dagger_core.h.

o21_im = 0

Definition at line 146 of file dw_dslash_dagger_core.h.

o21_re = 0

Definition at line 145 of file dw_dslash_dagger_core.h.

o22_im = 0

Definition at line 148 of file dw_dslash_dagger_core.h.

o22_re = 0

Definition at line 147 of file dw_dslash_dagger_core.h.

o30_im = 0

Definition at line 150 of file dw_dslash_dagger_core.h.

o30_re = 0

Definition at line 149 of file dw_dslash_dagger_core.h.

o31_im = 0

Definition at line 152 of file dw_dslash_dagger_core.h.

o31_re = 0

Definition at line 151 of file dw_dslash_dagger_core.h.

o32_im = 0

Definition at line 154 of file dw_dslash_dagger_core.h.

o32_re = 0

Definition at line 153 of file dw_dslash_dagger_core.h.

s_parity = ( sid/(X4*X3*X2*X1h) ) % 2

Definition at line 182 of file dw_dslash_dagger_core.h.

int sid = ((blockIdx.y*blockDim.y + threadIdx.y)*gridDim.x + blockIdx.x)*blockDim.x + threadIdx.x

Definition at line 177 of file dw_dslash_dagger_core.h.

X = 2*sid + (boundaryCrossing + param.parity) % 2

Definition at line 180 of file dw_dslash_dagger_core.h.

x1 = X % X1

Definition at line 180 of file dw_dslash_dagger_core.h.

x2 = (X/X1) % X2

Definition at line 180 of file dw_dslash_dagger_core.h.

x3 = (X/(X1*X2)) % X3

Definition at line 180 of file dw_dslash_dagger_core.h.

x4 = (X/(X1*X2*X3)) % X4

Definition at line 180 of file dw_dslash_dagger_core.h.

xs = X/(X1*X2*X3*X4)

Definition at line 180 of file dw_dslash_dagger_core.h.