Moved texture and constant definitions into header files

git-svn-id: http://lattice.bu.edu/qcdalg/cuda/quda@530 be54200a-260c-0410-bdd7-ce6af2a381ab

Makefile.tmpl

@@ -37,7 +37,7 @@ LDFLAGS = -fPIC $(LIB)
 all: dslash_test invert_test su3_test pack_test blas_test
 
 QUDA = libquda.a
-QUDA_OBJS = blas_quda.o blas_reference.o clover_field.o clover_quda.o dslash_quda.o      \
+QUDA_OBJS = blas_quda.o blas_reference.o clover_quda.o dslash_quda.o      \
 	dslash_reference.o gauge_quda.o inv_bicgstab_quda.o inv_cg_quda.o \
 	invert_quda.o spinor_quda.o util_quda.o
 QUDA_HDRS = blas_quda.h blas_reference.h clover_def.h dslash_def.h        \

dslash_constants.h

+__constant__ int X1h;
+__constant__ int X1;
+__constant__ int X2;
+__constant__ int X3;
+__constant__ int X4;
+
+__constant__ int X1m1;
+__constant__ int X2m1;
+__constant__ int X3m1;
+__constant__ int X4m1;
+
+__constant__ int X2X1mX1;
+__constant__ int X3X2X1mX2X1;
+__constant__ int X4X3X2X1mX3X2X1;
+__constant__ int X4X3X2X1hmX3X2X1h;
+
+__constant__ float X1h_inv;
+__constant__ float X2_inv;
+__constant__ float X3_inv;
+
+__constant__ int X2X1;
+__constant__ int X3X2X1;
+
+__constant__ int Vh;
+__constant__ int sp_stride;
+__constant__ int ga_stride;
+__constant__ int cl_stride;
+
+__constant__ int gauge_fixed;
+
+// single precision constants
+__constant__ float anisotropy_f;
+__constant__ float t_boundary_f;
+__constant__ float pi_f;
+
+// double precision constants
+__constant__ double anisotropy;
+__constant__ double t_boundary;
+

dslash_textures.h

+#if (__CUDA_ARCH__ == 130)
+static __inline__ __device__ double2 fetch_double2(texture<int4, 1> t, int i)
+{
+  int4 v = tex1Dfetch(t,i);
+  return make_double2(__hiloint2double(v.y, v.x), __hiloint2double(v.w, v.z));
+}
+#endif
+
+// Double precision gauge field
+texture<int4, 1> gauge0TexDouble;
+texture<int4, 1> gauge1TexDouble;
+
+// Single precision gauge field
+texture<float4, 1, cudaReadModeElementType> gauge0TexSingle;
+texture<float4, 1, cudaReadModeElementType> gauge1TexSingle;
+
+// Half precision gauge field
+texture<short4, 1, cudaReadModeNormalizedFloat> gauge0TexHalf;
+texture<short4, 1, cudaReadModeNormalizedFloat> gauge1TexHalf;
+
+// Double precision input spinor field
+texture<int4, 1> spinorTexDouble;
+
+// Single precision input spinor field
+texture<float4, 1, cudaReadModeElementType> spinorTexSingle;
+
+// Half precision input spinor field
+texture<short4, 1, cudaReadModeNormalizedFloat> spinorTexHalf;
+texture<float, 1, cudaReadModeElementType> spinorTexNorm;
+
+// Double precision accumulate spinor field
+texture<int4, 1> accumTexDouble;
+
+// Single precision accumulate spinor field
+texture<float4, 1, cudaReadModeElementType> accumTexSingle;
+
+// Half precision accumulate spinor field
+texture<short4, 1, cudaReadModeNormalizedFloat> accumTexHalf;
+texture<float, 1, cudaReadModeElementType> accumTexNorm;
+
+// Double precision clover term
+texture<int4, 1> cloverTexDouble;
+
+// Single precision clover term
+texture<float4, 1, cudaReadModeElementType> cloverTexSingle;
+
+// Half precision clover term
+texture<short4, 1, cudaReadModeNormalizedFloat> cloverTexHalf;
+texture<float, 1, cudaReadModeElementType> cloverTexNorm;