#include "read_gauge.h"
#include "io_spinor.h"

Macros
#define	DSLASH_SHARED_FLOATS_PER_THREAD 19

#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	acc00_re accum0.x

#define	acc00_im accum0.y

#define	acc01_re accum0.z

#define	acc01_im accum0.w

#define	acc02_re accum1.x

#define	acc02_im accum1.y

#define	acc10_re accum1.z

#define	acc10_im accum1.w

#define	acc11_re accum2.x

#define	acc11_im accum2.y

#define	acc12_re accum2.z

#define	acc12_im accum2.w

#define	acc20_re accum3.x

#define	acc20_im accum3.y

#define	acc21_re accum3.z

#define	acc21_im accum3.w

#define	acc22_re accum4.x

#define	acc22_im accum4.y

#define	acc30_re accum4.z

#define	acc30_im accum4.w

#define	acc31_re accum5.x

#define	acc31_im accum5.y

#define	acc32_re accum5.z

#define	acc32_im accum5.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	o00_re s[0*SHARED_STRIDE]

#define	o00_im s[1*SHARED_STRIDE]

#define	o01_re s[2*SHARED_STRIDE]

#define	o01_im s[3*SHARED_STRIDE]

#define	o02_re s[4*SHARED_STRIDE]

#define	o02_im s[5*SHARED_STRIDE]

#define	o10_re s[6*SHARED_STRIDE]

#define	o10_im s[7*SHARED_STRIDE]

#define	o11_re s[8*SHARED_STRIDE]

#define	o11_im s[9*SHARED_STRIDE]

#define	o12_re s[10*SHARED_STRIDE]

#define	o12_im s[11*SHARED_STRIDE]

#define	o20_re s[12*SHARED_STRIDE]

#define	o20_im s[13*SHARED_STRIDE]

#define	o21_re s[14*SHARED_STRIDE]

#define	o21_im s[15*SHARED_STRIDE]

#define	o22_re s[16*SHARED_STRIDE]

#define	o22_im s[17*SHARED_STRIDE]

#define	o30_re s[18*SHARED_STRIDE]

#define	SHARED_STRIDE 16

Functions
	if (sid >=param.threads) return

	coordsFromIndex< EVEN_X > (X, x1, x2, x3, x4, sid, param.parity, dims)

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

	READ_GAUGE_MATRIX (G, GAUGE0TEX, 0, ga_idx, ga_stride)

	RECONSTRUCT_GAUGE_MATRIX (0)

	READ_GAUGE_MATRIX (G, GAUGE1TEX, 1, ga_idx, ga_stride)

	RECONSTRUCT_GAUGE_MATRIX (1)

	READ_GAUGE_MATRIX (G, GAUGE0TEX, 2, ga_idx, ga_stride)

	RECONSTRUCT_GAUGE_MATRIX (2)

	READ_GAUGE_MATRIX (G, GAUGE1TEX, 3, ga_idx, ga_stride)

	RECONSTRUCT_GAUGE_MATRIX (3)

	READ_GAUGE_MATRIX (G, GAUGE0TEX, 4, ga_idx, ga_stride)

	RECONSTRUCT_GAUGE_MATRIX (4)

	READ_GAUGE_MATRIX (G, GAUGE1TEX, 5, ga_idx, ga_stride)

	RECONSTRUCT_GAUGE_MATRIX (5)

	if (gauge_fixed &&ga_idx< X4X3X2X1hmX3X2X1h)

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

	READ_GAUGE_MATRIX (G, GAUGE0TEX, 6, ga_idx, ga_stride)

	RECONSTRUCT_GAUGE_MATRIX (6)

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

	READ_GAUGE_MATRIX (G, GAUGE1TEX, 7, ga_idx, ga_stride)

	RECONSTRUCT_GAUGE_MATRIX (7)

	WRITE_SPINOR (param.sp_stride)

Variables
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

__shared__ char	s_data []

VOLATILE spinorFloat *	s

int	x1

int	x2

int	x3

int	x4

int	X

int	sid = blockIdx.x*blockDim.x + threadIdx.x

const int	dims [] = {X1, X2, X3, X4}

	o00_re = 0

	o00_im = 0

	o01_re = 0

	o01_im = 0

	o02_re = 0

	o02_im = 0

	o10_re = 0

	o10_im = 0

	o11_re = 0

	o11_im = 0

	o12_re = 0

	o12_im = 0

	o20_re = 0

	o20_im = 0

	o21_re = 0

	o21_im = 0

	o22_re = 0

	o22_im = 0

	o30_re = 0

const int	ga_idx = sid

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

	else