#include <stdlib.h>
#include <stdio.h>
#include <quda_internal.h>
#include <blas_quda.h>
#include <color_spinor_field.h>
#include <cuComplex.h>
#include <blas_param.h>
#include "reduce_core.h"
#include "reduce_complex_core.h"
#include "reduce_triple_core.h"

Defines
#define	REDUCE_MAX_BLOCKS 65536
#define	REDUCE_DOUBLE 64
#define	REDUCE_KAHAN 32
#define	REDUCE_TYPE REDUCE_KAHAN
#define	QudaSumFloat float
#define	QudaSumComplex cuComplex
#define	QudaSumFloat3 float3
#define	EMUSYNC
#define	READ_DOUBLE2_TEXTURE(x, i) fetch_double2(x##TexDouble2, i)
#define	READ_FLOAT2_TEXTURE(x, i) tex1Dfetch(x##TexSingle2, i)
#define	RECONSTRUCT_HALF_SPINOR(a, texHalf, texNorm, length)
#define	RECONSTRUCT_HALF_SPINOR_ST(a, texHalf, texNorm, length)
#define	READ_HALF_SPINOR_TEX(a, tex, texNorm, length)
#define	READ_HALF_SPINOR(a, tex, length)
#define	READ_HALF_SPINOR_ST(a, tex, length)
#define	SHORT_LENGTH 65536
#define	SCALE_FLOAT ((SHORT_LENGTH-1) * 0.5)
#define	SHIFT_FLOAT (-1.f / (SHORT_LENGTH-1))
#define	FAST_ABS_MAX(a, b) fmaxf(fabsf(a), fabsf(b));
#define	FAST_MAX(a, b) fmaxf(a, b);
#define	CONSTRUCT_HALF_SPINOR_FROM_SINGLE(h, n, a, length)
#define	CONSTRUCT_HALF_SPINOR_FROM_DOUBLE(h, n, a, length)
#define	CONSTRUCT_HALF_SPINOR_FROM_SINGLE_ST(h, n, a, length)
#define	CONSTRUCT_HALF_SPINOR_FROM_DOUBLE_ST(h, n, a, length)
#define	SUM_FLOAT4(sum, a) float sum = a.x + a.y + a.z + a.w;
#define	SUM_FLOAT2(sum, a) float sum = a.x + a.y;
#define	REAL_DOT_FLOAT4(dot, a, b) float dot = a.xb.x + a.yb.y + a.zb.z + a.wb.w;
#define	REAL_DOT_FLOAT2(dot, a, b) float dot = a.xb.x + a.yb.y;
#define	IMAG_DOT_FLOAT4(dot, a, b) float dot = a.xb.y - a.yb.x + a.zb.w - a.wb.z;
#define	IMAG_DOT_FLOAT2(dot, a, b) float dot = a.xb.y - a.yb.x;
#define	AX_FLOAT4(a, X) X.x = a; X.y = a; X.z = a; X.w = a;
#define	AX_FLOAT2(a, X) X.x = a; X.y = a;
#define	XPY_FLOAT4(X, Y) Y.x += X.x; Y.y += X.y; Y.z += X.z; Y.w += X.w;
#define	XPY_FLOAT2(X, Y) Y.x += X.x; Y.y += X.y;
#define	XMY_FLOAT4(X, Y) Y.x = X.x - Y.x; Y.y = X.y - Y.y; Y.z = X.z - Y.z; Y.w = X.w - Y.w;
#define	XMY_FLOAT2(X, Y) Y.x = X.x - Y.x; Y.y = X.y - Y.y;
#define	MXPY_FLOAT4(X, Y) Y.x -= X.x; Y.y -= X.y; Y.z -= X.z; Y.w -= X.w;
#define	MXPY_FLOAT2(X, Y) Y.x -= X.x; Y.y -= X.y;
#define	AXPY_FLOAT4(a, X, Y)
#define	AXPY_FLOAT2(a, X, Y) Y.x += aX.x; Y.y += aX.y;
#define	AXPBY_FLOAT4(a, X, b, Y)
#define	AXPBY_FLOAT2(a, X, b, Y) Y.x = bY.x; Y.x += aX.x; Y.y = bY.y; Y.y += aX.y; \
#define	XPAY_FLOAT4(X, a, Y)
#define	XPAY_FLOAT2(X, a, Y) Y.x = X.x + aY.x; Y.y = X.y + aY.y;
#define	CAXPY_FLOAT4(a, X, Y)
#define	CAXPY_FLOAT2(a, X, Y)
#define	CMAXPY_FLOAT4(a, X, Y)
#define	CAXPBY_FLOAT4(a, X, b, Y)
#define	CAXPBY_FLOAT2(a, X, b, Y)
#define	CXPAYPBZ_FLOAT4(X, a, Y, b, Z)
#define	CXPAYPBZ_FLOAT2(X, a, Y, b, Z)
#define	CAXPBYPZ_FLOAT4(a, X, b, Y, Z)
#define	CAXPBYPZ_FLOAT2(a, X, b, Y, Z)
#define	checkSpinor(a, b)
#define	checkSpinorLength(a, b)
#define	REDUCE_FUNC_NAME(suffix) sumD##suffix
#define	REDUCE_TYPES Float *a
#define	REDUCE_PARAMS a
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) a[i]
#define	REDUCE_FUNC_NAME(suffix) sumS##suffix
#define	REDUCE_TYPES Float *a
#define	REDUCE_PARAMS a
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) a[i].x + a[i].y
#define	REDUCE_FUNC_NAME(suffix) sumH##suffix
#define	REDUCE_TYPES Float *aN, int stride
#define	REDUCE_PARAMS aN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (ac*s0)
#define	REDUCE_FUNC_NAME(suffix) sumHSt##suffix
#define	REDUCE_TYPES Float *aN, int stride
#define	REDUCE_PARAMS aN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (ac*s0)
#define	REDUCE_FUNC_NAME(suffix) normD##suffix
#define	REDUCE_TYPES Float *a
#define	REDUCE_PARAMS a
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (a[i]*a[i])
#define	REDUCE_FUNC_NAME(suffix) normS##suffix
#define	REDUCE_TYPES Float *a
#define	REDUCE_PARAMS a
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (a[i].xa[i].x + a[i].ya[i].y)
#define	REDUCE_FUNC_NAME(suffix) normH##suffix
#define	REDUCE_TYPES Float *aN, int stride
#define	REDUCE_PARAMS aN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (acacnorm0)
#define	REDUCE_FUNC_NAME(suffix) normHSt##suffix
#define	REDUCE_TYPES Float *aN, int stride
#define	REDUCE_PARAMS aN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (acacnorm0)
#define	REDUCE_FUNC_NAME(suffix) reDotProductD##suffix
#define	REDUCE_TYPES Float a, Float b
#define	REDUCE_PARAMS a, b
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (a[i]*b[i])
#define	REDUCE_FUNC_NAME(suffix) reDotProductS##suffix
#define	REDUCE_TYPES Float a, Float b
#define	REDUCE_PARAMS a, b
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (a[i].xb[i].x + a[i].yb[i].y)
#define	REDUCE_FUNC_NAME(suffix) reDotProductH##suffix
#define	REDUCE_TYPES Float aN, Float bN, int stride
#define	REDUCE_PARAMS aN, bN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (acbcrdot0)
#define	REDUCE_FUNC_NAME(suffix) reDotProductHSt##suffix
#define	REDUCE_TYPES Float aN, Float bN, int stride
#define	REDUCE_PARAMS aN, bN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (acbcrdot0)
#define	REDUCE_FUNC_NAME(suffix) axpyNormF##suffix
#define	REDUCE_TYPES Float a, Float x, Float y
#define	REDUCE_PARAMS a, x, y
#define	REDUCE_AUXILIARY(i) y[i] = a*x[i] + y[i]
#define	REDUCE_OPERATION(i) (y[i]*y[i])
#define	REDUCE_FUNC_NAME(suffix) axpyNormH##suffix
#define	REDUCE_TYPES Float a, short4 yH, float yN, int stride
#define	REDUCE_PARAMS a, yH, yN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (norm0)
#define	REDUCE_FUNC_NAME(suffix) axpyNormH##suffix
#define	REDUCE_TYPES Float a, short2 yH, float yN, int stride
#define	REDUCE_PARAMS a, yH, yN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (norm0)
#define	REDUCE_FUNC_NAME(suffix) xmyNormF##suffix
#define	REDUCE_TYPES Float x, Float y
#define	REDUCE_PARAMS x, y
#define	REDUCE_AUXILIARY(i) y[i] = x[i] - y[i]
#define	REDUCE_OPERATION(i) (y[i]*y[i])
#define	REDUCE_FUNC_NAME(suffix) xmyNormH##suffix
#define	REDUCE_TYPES Float d1, Float d2, short4 yH, float yN, int stride
#define	REDUCE_PARAMS d1, d2, yH, yN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (norm0)
#define	REDUCE_FUNC_NAME(suffix) xmyNormH##suffix
#define	REDUCE_TYPES Float d1, Float d2, short2 yH, float yN, int stride
#define	REDUCE_PARAMS d1, d2, yH, yN, stride
#define	REDUCE_AUXILIARY(i)
#define	REDUCE_OPERATION(i) (norm0)
#define	REDUCE_FUNC_NAME(suffix) cDotProductD##suffix
#define	REDUCE_TYPES Float2 x, Float2 y, Float c
#define	REDUCE_PARAMS x, y, c
#define	REDUCE_REAL_AUXILIARY(i) Float2 a = READ_DOUBLE2_TEXTURE(x, i);
#define	REDUCE_IMAG_AUXILIARY(i) Float2 b = READ_DOUBLE2_TEXTURE(y, i);
#define	REDUCE_REAL_OPERATION(i) (a.xb.x + a.yb.y)
#define	REDUCE_IMAG_OPERATION(i) (a.xb.y - a.yb.x)
#define	REDUCE_FUNC_NAME(suffix) cDotProductS##suffix
#define	REDUCE_TYPES Float2 x, Float2 y, Float c
#define	REDUCE_PARAMS x, y, c
#define	REDUCE_REAL_AUXILIARY(i) Float2 a = read_Float2(x, i);
#define	REDUCE_IMAG_AUXILIARY(i) Float2 b = read_Float2(y, i);
#define	REDUCE_REAL_OPERATION(i) (a.xb.x + a.yb.y)
#define	REDUCE_IMAG_OPERATION(i) (a.xb.y - a.yb.x)
#define	REDUCE_FUNC_NAME(suffix) cDotProductH##suffix
#define	REDUCE_TYPES Float aN, Float2 bN, int stride
#define	REDUCE_PARAMS aN, bN, stride
#define	REDUCE_REAL_AUXILIARY(i)
#define	REDUCE_IMAG_AUXILIARY(i)
#define	REDUCE_REAL_OPERATION(i) (acbcrdot0)
#define	REDUCE_IMAG_OPERATION(i) (acbcidot0)
#define	REDUCE_FUNC_NAME(suffix) cDotProductHSt##suffix
#define	REDUCE_TYPES Float aN, Float2 bN, int stride
#define	REDUCE_PARAMS aN, bN, stride
#define	REDUCE_REAL_AUXILIARY(i)
#define	REDUCE_IMAG_AUXILIARY(i)
#define	REDUCE_REAL_OPERATION(i) (acbcrdot0)
#define	REDUCE_IMAG_OPERATION(i) (acbcidot0)
#define	REDUCE_FUNC_NAME(suffix) xpaycDotzyD##suffix
#define	REDUCE_TYPES Float2 x, Float a, Float2 y, Float2 *z
#define	REDUCE_PARAMS x, a, y, z
#define	REDUCE_REAL_AUXILIARY(i)
#define	REDUCE_IMAG_AUXILIARY(i) y[i].x = X.x + aY.x; y[i].y = X.y + aY.y
#define	REDUCE_REAL_OPERATION(i) (Z.xy[i].x + Z.yy[i].y)
#define	REDUCE_IMAG_OPERATION(i) (Z.xy[i].y - Z.yy[i].x)
#define	REDUCE_FUNC_NAME(suffix) xpaycDotzyS##suffix
#define	REDUCE_TYPES Float2 x, Float a, Float2 y, Float2 *z
#define	REDUCE_PARAMS x, a, y, z
#define	REDUCE_REAL_AUXILIARY(i) y[i].x = x[i].x + a*y[i].x
#define	REDUCE_IMAG_AUXILIARY(i) y[i].y = x[i].y + a*y[i].y
#define	REDUCE_REAL_OPERATION(i) (z[i].xy[i].x + z[i].yy[i].y)
#define	REDUCE_IMAG_OPERATION(i) (z[i].xy[i].y - z[i].yy[i].x)
#define	REDUCE_FUNC_NAME(suffix) xpaycDotzyH##suffix
#define	REDUCE_TYPES Float a, short4 yH, Float2 yN, int stride
#define	REDUCE_PARAMS a, yH, yN, stride
#define	REDUCE_REAL_AUXILIARY(i)
#define	REDUCE_IMAG_AUXILIARY(i)
#define	REDUCE_REAL_OPERATION(i) (rdot0)
#define	REDUCE_IMAG_OPERATION(i) (idot0)
#define	REDUCE_FUNC_NAME(suffix) xpaycDotzyH##suffix
#define	REDUCE_TYPES Float a, short2 yH, Float2 yN, int stride
#define	REDUCE_PARAMS a, yH, yN, stride
#define	REDUCE_REAL_AUXILIARY(i)
#define	REDUCE_IMAG_AUXILIARY(i)
#define	REDUCE_REAL_OPERATION(i) (rdot0)
#define	REDUCE_IMAG_OPERATION(i) (idot0)
#define	REDUCE_FUNC_NAME(suffix) cDotProductNormAD##suffix
#define	REDUCE_TYPES Float2 x, Float2 y
#define	REDUCE_PARAMS x, y
#define	REDUCE_X_AUXILIARY(i) Float2 a = READ_DOUBLE2_TEXTURE(x, i);
#define	REDUCE_Y_AUXILIARY(i) Float2 b = READ_DOUBLE2_TEXTURE(y, i);
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (a.xb.x + a.yb.y)
#define	REDUCE_Y_OPERATION(i) (a.xb.y - a.yb.x)
#define	REDUCE_Z_OPERATION(i) (a.xa.x + a.ya.y)
#define	REDUCE_FUNC_NAME(suffix) cDotProductNormAS##suffix
#define	REDUCE_TYPES Float2 a, Float2 b
#define	REDUCE_PARAMS a, b
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (a[i].xb[i].x + a[i].yb[i].y)
#define	REDUCE_Y_OPERATION(i) (a[i].xb[i].y - a[i].yb[i].x)
#define	REDUCE_Z_OPERATION(i) (a[i].xa[i].x + a[i].ya[i].y)
#define	REDUCE_FUNC_NAME(suffix) cDotProductNormAH##suffix
#define	REDUCE_TYPES Float2 xN, Float2 yN, int stride
#define	REDUCE_PARAMS xN, yN, stride
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (xcycrdot0)
#define	REDUCE_Y_OPERATION(i) (xcycidot0)
#define	REDUCE_Z_OPERATION(i) (xcxcnorm0)
#define	REDUCE_FUNC_NAME(suffix) cDotProductNormAHSt##suffix
#define	REDUCE_TYPES Float2 xN, Float2 yN, int stride
#define	REDUCE_PARAMS xN, yN, stride
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (xcycrdot0)
#define	REDUCE_Y_OPERATION(i) (xcycidot0)
#define	REDUCE_Z_OPERATION(i) (xcxcnorm0)
#define	REDUCE_FUNC_NAME(suffix) cDotProductNormBD##suffix
#define	REDUCE_TYPES Float2 x, Float2 y
#define	REDUCE_PARAMS x, y
#define	REDUCE_X_AUXILIARY(i) Float2 a = READ_DOUBLE2_TEXTURE(x, i);
#define	REDUCE_Y_AUXILIARY(i) Float2 b = READ_DOUBLE2_TEXTURE(y, i);
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (a.xb.x + a.yb.y)
#define	REDUCE_Y_OPERATION(i) (a.xb.y - a.yb.x)
#define	REDUCE_Z_OPERATION(i) (b.xb.x + b.yb.y)
#define	REDUCE_FUNC_NAME(suffix) cDotProductNormBS##suffix
#define	REDUCE_TYPES Float2 a, Float2 b
#define	REDUCE_PARAMS a, b
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (a[i].xb[i].x + a[i].yb[i].y)
#define	REDUCE_Y_OPERATION(i) (a[i].xb[i].y - a[i].yb[i].x)
#define	REDUCE_Z_OPERATION(i) (b[i].xb[i].x + b[i].yb[i].y)
#define	REDUCE_FUNC_NAME(suffix) cDotProductNormBH##suffix
#define	REDUCE_TYPES Float2 xN, Float2 yN, int stride
#define	REDUCE_PARAMS xN, yN, stride
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (xcycrdot0)
#define	REDUCE_Y_OPERATION(i) (xcycidot0)
#define	REDUCE_Z_OPERATION(i) (ycycnorm0)
#define	REDUCE_FUNC_NAME(suffix) cDotProductNormBHSt##suffix
#define	REDUCE_TYPES Float2 xN, Float2 yN, int stride
#define	REDUCE_PARAMS xN, yN, stride
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (xcycrdot0)
#define	REDUCE_Y_OPERATION(i) (xcycidot0)
#define	REDUCE_Z_OPERATION(i) (ycycnorm0)
#define	REDUCE_FUNC_NAME(suffix) caxpbypzYmbwcDotProductWYNormYD##suffix
#define	REDUCE_TYPES Float2 a, Float2 x, Float2 b, Float2 y, Float2 z, Float2 w, Float2 *u
#define	REDUCE_PARAMS a, x, b, y, z, w, u
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (u[i].xy[i].x + u[i].yy[i].y)
#define	REDUCE_Y_OPERATION(i) (u[i].xy[i].y - u[i].yy[i].x)
#define	REDUCE_Z_OPERATION(i) (y[i].xy[i].x + y[i].yy[i].y)
#define	REDUCE_FUNC_NAME(suffix) caxpbypzYmbwcDotProductWYNormYS##suffix
#define	REDUCE_TYPES Float2 a, Float2 x, Float2 b, Float2 y, Float2 z, Float2 w, Float2 *u
#define	REDUCE_PARAMS a, x, b, y, z, w, u
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (u[i].xy[i].x + u[i].yy[i].y)
#define	REDUCE_Y_OPERATION(i) (u[i].xy[i].y - u[i].yy[i].x)
#define	REDUCE_Z_OPERATION(i) (y[i].xy[i].x + y[i].yy[i].y)
#define	REDUCE_FUNC_NAME(suffix) caxpbypzYmbwcDotProductWYNormYH##suffix
#define	REDUCE_TYPES Float2 a, Float2 b, short4 yH, float yN, short4 zH, float zN, float wN, float uN, int stride
#define	REDUCE_PARAMS a, b, yH, yN, zH, zN, wN, uN, stride
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (uc*rdot0)
#define	REDUCE_Y_OPERATION(i) (uc*idot0)
#define	REDUCE_Z_OPERATION(i) (norm0)
#define	REDUCE_FUNC_NAME(suffix) caxpbypzYmbwcDotProductWYNormYH##suffix
#define	REDUCE_TYPES Float2 a, Float2 b, short2 yH, float yN, short2 zH, float zN, float wN, float uN, int stride
#define	REDUCE_PARAMS a, b, yH, yN, zH, zN, wN, uN, stride
#define	REDUCE_X_AUXILIARY(i)
#define	REDUCE_Y_AUXILIARY(i)
#define	REDUCE_Z_AUXILIARY(i)
#define	REDUCE_X_OPERATION(i) (uc*rdot0)
#define	REDUCE_Y_OPERATION(i) (uc*idot0)
#define	REDUCE_Z_OPERATION(i) (norm0)
Functions
double2	operator+ (const double2 &x, const double2 &y)
double3	operator+ (const double3 &x, const double3 &y)
__device__ float2	operator* (const float a, const float2 x)
template<typename Float2 >
__device__ Float2	operator+ (const Float2 x, const Float2 y)
template<typename Float2 >
__device__ Float2	operator+= (Float2 &x, const Float2 y)
template<typename Float2 >
__device__ Float2	operator-= (Float2 &x, const Float2 y)
template<typename Float , typename Float2 >
__device__ Float2	operator*= (Float2 &x, const Float a)
template<typename Float >
__device__ float4	operator*= (float4 &a, const Float &b)
void	zeroCuda (cudaColorSpinorField &a)
void	initBlas (void)
void	endBlas (void)
void	setBlasTuning (int tuning)
void	setBlasParam (int kernel, int prec, int threads, int blocks)
void	setBlock (int kernel, int length, QudaPrecision precision)
float2 __device__	read_Float2 (float2 *x, int i)
double2 __device__	read_Float2 (double2 *x, int i)
float2 __device__	make_Float2 (float2 x)
double2 __device__	make_Float2 (double2 x)
__device__ short	float2short (float c, float a)
__device__ float	short2float (short a)
__device__ short4	float42short4 (float c, float4 a)
__device__ float	fast_abs_max (float4 a)
__global__ void	convertDSKernel (double2 dst, float4 src, int length)
__global__ void	convertDSKernel (double2 dst, float2 src, int length)
__global__ void	convertSDKernel (float4 dst, double2 src, int length)
__global__ void	convertSDKernel (float2 dst, double2 src, int length)
__global__ void	convertHSKernel (short4 h, float norm, int length, int real_length)
__global__ void	convertHSKernel (short2 h, float norm, int length, int real_length)
__global__ void	convertSHKernel (float4 *res, int length, int real_length)
__global__ void	convertSHKernel (float2 *res, int length, int real_length)
__global__ void	convertHDKernel (short4 h, float norm, int length, int real_length)
__global__ void	convertHDKernel (short2 h, float norm, int length, int real_length)
__global__ void	convertDHKernel (double2 *res, int length, int real_length)
__global__ void	convertDHKernelSt (double2 *res, int length, int real_length)
void	copyCuda (cudaColorSpinorField &dst, const cudaColorSpinorField &src)
template<typename Float , typename Float2 >
__global__ void	axpbyKernel (Float a, Float2 x, Float b, Float2 y, int length)
__global__ void	axpbyHKernel (float a, float b, short4 yH, float yN, int stride, int length)
__global__ void	axpbyHKernel (float a, float b, short2 yH, float yN, int stride, int length)
void	axpbyCuda (const double &a, cudaColorSpinorField &x, const double &b, cudaColorSpinorField &y)
template<typename Float >
__global__ void	xpyKernel (Float x, Float y, int len)
__global__ void	xpyHKernel (short4 yH, float yN, int stride, int length)
__global__ void	xpyHKernel (short2 yH, float yN, int stride, int length)
void	xpyCuda (cudaColorSpinorField &x, cudaColorSpinorField &y)
template<typename Float , typename Float2 >
__global__ void	axpyKernel (Float a, Float2 x, Float2 y, int len)
__global__ void	axpyHKernel (float a, short4 yH, float yN, int stride, int length)
__global__ void	axpyHKernel (float a, short2 yH, float yN, int stride, int length)
void	axpyCuda (const double &a, cudaColorSpinorField &x, cudaColorSpinorField &y)
template<typename Float , typename Float2 >
__global__ void	xpayKernel (const Float2 x, Float a, Float2 y, int len)
__global__ void	xpayHKernel (float a, short4 yH, float yN, int stride, int length)
__global__ void	xpayHKernel (float a, short2 yH, float yN, int stride, int length)
void	xpayCuda (const cudaColorSpinorField &x, const double &a, cudaColorSpinorField &y)
template<typename Float >
__global__ void	mxpyKernel (Float x, Float y, int len)
__global__ void	mxpyHKernel (short4 yH, float yN, int stride, int length)
__global__ void	mxpyHKernel (short2 yH, float yN, int stride, int length)
void	mxpyCuda (cudaColorSpinorField &x, cudaColorSpinorField &y)
template<typename Float , typename Float2 >
__global__ void	axKernel (Float a, Float2 *x, int len)
__global__ void	axHKernel (float a, short4 xH, float xN, int stride, int length)
__global__ void	axHKernel (float a, short2 xH, float xN, int stride, int length)
void	axCuda (const double &a, cudaColorSpinorField &x)
template<typename Float2 >
__global__ void	caxpyDKernel (Float2 a, Float2 x, Float2 y, int len)
template<typename Float2 >
__global__ void	caxpySKernel (Float2 a, Float2 x, Float2 y, int len)
__global__ void	caxpyHKernel (float2 a, short4 yH, float yN, int stride, int length)
__global__ void	caxpyHKernel (float2 a, short2 yH, float yN, int stride, int length)
void	caxpyCuda (const Complex &a, cudaColorSpinorField &x, cudaColorSpinorField &y)
template<typename Float2 >
__global__ void	caxpbyDKernel (Float2 a, Float2 x, Float2 b, Float2 y, int len)
template<typename Float2 >
__global__ void	caxpbySKernel (Float2 a, Float2 x, Float2 b, Float2 y, int len)
__global__ void	caxpbyHKernel (float2 a, float2 b, short4 yH, float yN, int stride, int length)
__global__ void	caxpbyHKernel (float2 a, float2 b, short2 yH, float yN, int stride, int length)
void	caxpbyCuda (const Complex &a, cudaColorSpinorField &x, const Complex &b, cudaColorSpinorField &y)
template<typename Float2 >
__global__ void	cxpaypbzDKernel (Float2 x, Float2 a, Float2 y, Float2 b, Float2 *z, int len)
template<typename Float2 >
__global__ void	cxpaypbzSKernel (Float2 x, Float2 a, Float2 y, Float2 b, Float2 *z, int len)
__global__ void	cxpaypbzHKernel (float2 a, float2 b, short4 zH, float zN, int stride, int length)
__global__ void	cxpaypbzHKernel (float2 a, float2 b, short2 zH, float zN, int stride, int length)
void	cxpaypbzCuda (cudaColorSpinorField &x, const Complex &a, cudaColorSpinorField &y, const Complex &b, cudaColorSpinorField &z)
template<typename Float , typename Float2 >
__global__ void	axpyBzpcxDKernel (Float a, Float2 x, Float2 y, Float b, Float2 *z, Float c, int len)
template<typename Float , typename Float2 >
__global__ void	axpyBzpcxSKernel (Float a, Float2 x, Float2 y, Float b, Float2 *z, Float c, int len)
__global__ void	axpyBzpcxHKernel (float a, float b, float c, short4 xH, float xN, short4 yH, float yN, int stride, int length)
__global__ void	axpyBzpcxHKernel (float a, float b, float c, short2 xH, float xN, short2 yH, float yN, int stride, int length)
void	axpyBzpcxCuda (const double &a, cudaColorSpinorField &x, cudaColorSpinorField &y, const double &b, cudaColorSpinorField &z, const double &c)
template<typename Float , typename Float2 >
__global__ void	axpyZpbxDKernel (Float a, Float2 x, Float2 y, Float2 *z, Float b, int len)
template<typename Float , typename Float2 >
__global__ void	axpyZpbxSKernel (Float a, Float2 x, Float2 y, Float2 *z, Float b, int len)
__global__ void	axpyZpbxHKernel (float a, float b, short4 xH, float xN, short4 yH, float yN, int stride, int length)
__global__ void	axpyZpbxHKernel (float a, float b, short2 xH, float xN, short2 yH, float yN, int stride, int length)
void	axpyZpbxCuda (const double &a, cudaColorSpinorField &x, cudaColorSpinorField &y, cudaColorSpinorField &z, const double &b)
template<typename Float2 >
__global__ void	caxpbypzYmbwDKernel (Float2 a, Float2 x, Float2 b, Float2 y, Float2 z, Float2 w, int len)
template<typename Float2 >
__global__ void	caxpbypzYmbwSKernel (Float2 a, Float2 x, Float2 b, Float2 y, Float2 z, Float2 w, int len)
__global__ void	caxpbypzYmbwHKernel (float2 a, float2 b, float xN, short4 yH, float yN, short4 zH, float zN, float wN, int stride, int length)
__global__ void	caxpbypzYmbwHKernel (float2 a, float2 b, float xN, short2 yH, float yN, short2 zH, float zN, float wN, int stride, int length)
void	caxpbypzYmbwCuda (const Complex &a, cudaColorSpinorField &x, const Complex &b, cudaColorSpinorField &y, cudaColorSpinorField &z, cudaColorSpinorField &w)
__device__ void	dsadd (volatile QudaSumFloat &c0, volatile QudaSumFloat &c1, const volatile QudaSumFloat &a0, const volatile QudaSumFloat &a1, const float b0, const float b1)
__device__ void	zcadd (volatile QudaSumComplex &c0, volatile QudaSumComplex &c1, const volatile QudaSumComplex &a0, const volatile QudaSumComplex &a1, const volatile QudaSumComplex &b0, const volatile QudaSumComplex &b1)
__device__ void	dsadd3 (volatile QudaSumFloat3 &c0, volatile QudaSumFloat3 &c1, const volatile QudaSumFloat3 &a0, const volatile QudaSumFloat3 &a1, const volatile QudaSumFloat3 &b0, const volatile QudaSumFloat3 &b1)
double	sumCuda (cudaColorSpinorField &a)
double	normCuda (const cudaColorSpinorField &a)
double	reDotProductCuda (cudaColorSpinorField &a, cudaColorSpinorField &b)
double	axpyNormCuda (const double &a, cudaColorSpinorField &x, cudaColorSpinorField &y)
double	xmyNormCuda (cudaColorSpinorField &x, cudaColorSpinorField &y)
Complex	cDotProductCuda (cudaColorSpinorField &x, cudaColorSpinorField &y)
Complex	xpaycDotzyCuda (cudaColorSpinorField &x, const double &a, cudaColorSpinorField &y, cudaColorSpinorField &z)
double3	cDotProductNormACuda (cudaColorSpinorField &x, cudaColorSpinorField &y)
double3	cDotProductNormBCuda (cudaColorSpinorField &x, cudaColorSpinorField &y)
double3	caxpbypzYmbwcDotProductWYNormYCuda (const Complex &a, cudaColorSpinorField &x, const Complex &b, cudaColorSpinorField &y, cudaColorSpinorField &z, cudaColorSpinorField &w, cudaColorSpinorField &u)
Variables
unsigned long long	blas_quda_flops
unsigned long long	blas_quda_bytes
texture< int4, 1 >	xTexDouble2
texture< int4, 1 >	yTexDouble2
texture< int4, 1 >	zTexDouble2
texture< int4, 1 >	wTexDouble2
texture< int4, 1 >	uTexDouble2
texture< float2, 1 >	xTexSingle2
texture< float2, 1 >	yTexSingle2
texture< float4, 1 >	xTexSingle4
texture< short4, 1, cudaReadModeNormalizedFloat >	texHalf1
texture< short2, 1, cudaReadModeNormalizedFloat >	texHalfSt1
texture< float, 1, cudaReadModeElementType >	texNorm1
texture< short4, 1, cudaReadModeNormalizedFloat >	texHalf2
texture< short2, 1, cudaReadModeNormalizedFloat >	texHalfSt2
texture< float, 1, cudaReadModeElementType >	texNorm2
texture< short4, 1, cudaReadModeNormalizedFloat >	texHalf3
texture< short2, 1, cudaReadModeNormalizedFloat >	texHalfSt3
texture< float, 1, cudaReadModeElementType >	texNorm3
texture< short4, 1, cudaReadModeNormalizedFloat >	texHalf4
texture< short2, 1, cudaReadModeNormalizedFloat >	texHalfSt4
texture< float, 1, cudaReadModeElementType >	texNorm4
texture< short4, 1, cudaReadModeNormalizedFloat >	texHalf5
texture< short2, 1, cudaReadModeNormalizedFloat >	texHalfSt5
texture< float, 1, cudaReadModeElementType >	texNorm5

Define Documentation

#define AX_FLOAT2	(	a,
		X
	)	X.x = a; X.y = a;

Definition at line 438 of file blas_quda.cu.

#define AX_FLOAT4	(	a,
		X
	)	X.x = a; X.y = a; X.z = a; X.w = a;

Definition at line 435 of file blas_quda.cu.

#define AXPBY_FLOAT2	(	a,
		X,
		b,
		Y
	)	Y.x = bY.x; Y.x += aX.x; Y.y = bY.y; Y.y += aX.y; \

Definition at line 476 of file blas_quda.cu.

#define AXPBY_FLOAT4	(	a,
		X,
		b,
		Y
	)

Value:

Y.x = b*Y.x; Y.x += a*X.x; Y.y = b*Y.y; Y.y += a*X.y;           \
  Y.z = b*Y.z; Y.z += a*X.z; Y.w = b*Y.w; Y.w += a*X.w;

Definition at line 472 of file blas_quda.cu.

#define AXPY_FLOAT2	(	a,
		X,
		Y
	)	Y.x += aX.x; Y.y += aX.y;

Definition at line 469 of file blas_quda.cu.

#define AXPY_FLOAT4	(	a,
		X,
		Y
	)

Value:

Y.x += a*X.x;   Y.y += a*X.y;                \
  Y.z += a*X.z; Y.w += a*X.w;

Definition at line 460 of file blas_quda.cu.

#define CAXPBY_FLOAT2	(	a,
		X,
		b,
		Y
	)

Value:

{ float2 y;                                                             \
  y.x = a.x*X.x; y.x -= a.y*X.y; y.x += b.x*Y.x; y.x -= b.y*Y.y;        \
  y.y = a.y*X.x; y.y += a.x*X.y; y.y += b.y*Y.x; y.y += b.x*Y.y;        \
  Y.x = y.x; Y.y = y.y;}

Definition at line 531 of file blas_quda.cu.

#define CAXPBY_FLOAT4	(	a,
		X,
		b,
		Y
	)

Value:

{ float2 y;                                                             \
  y.x = a.x*X.x; y.x -= a.y*X.y; y.x += b.x*Y.x; y.x -= b.y*Y.y;        \
  y.y = a.y*X.x; y.y += a.x*X.y; y.y += b.y*Y.x; y.y += b.x*Y.y;        \
  Y.x = y.x; Y.y = y.y;                                                 \
  y.x = a.x*X.z; y.x -= a.y*X.w; y.x += b.x*Y.z; y.x -= b.y*Y.w;        \
  y.y = a.y*X.z; y.y += a.x*X.w; y.y += b.y*Y.z; y.y += b.x*Y.w;        \
  Y.z = y.x; Y.w = y.y;}

Definition at line 522 of file blas_quda.cu.

#define CAXPBYPZ_FLOAT2	(	a,
		X,
		b,
		Y,
		Z
	)

Value:

Z.x += a.x*X.x - a.y*X.y + b.x*Y.x - b.y*Y.y;   \
  Z.y += a.y*X.x + a.x*X.y + b.y*Y.x + b.x*Y.y;

Definition at line 568 of file blas_quda.cu.

#define CAXPBYPZ_FLOAT4	(	a,
		X,
		b,
		Y,
		Z
	)

Value:

Z.x += a.x*X.x - a.y*X.y + b.x*Y.x - b.y*Y.y;   \
  Z.y += a.y*X.x + a.x*X.y + b.y*Y.x + b.x*Y.y;   \
  Z.z += a.x*X.z - a.y*X.w + b.x*Y.z - b.y*Y.w;   \
  Z.w += a.y*X.z + a.x*X.w + b.y*Y.z + b.x*Y.w;

Definition at line 554 of file blas_quda.cu.

#define CAXPY_FLOAT2	(	a,
		X,
		Y
	)

Value:

Y.x += a.x*X.x; Y.x -= a.y*X.y; \
  Y.y += a.y*X.x; Y.y += a.x*X.y;

Definition at line 507 of file blas_quda.cu.

#define CAXPY_FLOAT4	(	a,
		X,
		Y
	)

Value:

Y.x += a.x*X.x; Y.x -= a.y*X.y; \
  Y.y += a.y*X.x; Y.y += a.x*X.y;       \
  Y.z += a.x*X.z; Y.z -= a.y*X.w;       \
  Y.w += a.y*X.z; Y.w += a.x*X.w;

Definition at line 493 of file blas_quda.cu.

#define checkSpinor	(	a,
		b
	)

Value:

{                                                                       \
    if (a.Precision() != b.Precision())                                 \
      errorQuda("precisions do not match: %d %d", a.Precision(), b.Precision()); \
    if (a.Length() != b.Length())                                       \
      errorQuda("lengths do not match: %d %d", a.Length(), b.Length()); \
    if (a.Stride() != b.Stride())                                       \
      errorQuda("strides do not match: %d %d", a.Stride(), b.Stride()); \
  }

Definition at line 615 of file blas_quda.cu.

#define checkSpinorLength	(	a,
		b
	)

Value:

{                                                                       \
    if (a.Length() != b.Length()) {                                     \
      errorQuda("engths do not match: %d %d", a.Length(), b.Length());  \
    }

Definition at line 626 of file blas_quda.cu.

#define CMAXPY_FLOAT4	(	a,
		X,
		Y
	)

Value:

Y.x -= a.x*X.x; Y.x += a.y*X.y;         \
  Y.y -= a.y*X.x; Y.y -= a.x*X.y;               \
  Y.z -= a.x*X.z; Y.z += a.y*X.w;               \
  Y.w -= a.y*X.z; Y.w -= a.x*X.w;

Definition at line 516 of file blas_quda.cu.

#define CONSTRUCT_HALF_SPINOR_FROM_DOUBLE	(	h,
		n,
		a,
		length
	)

Value:

{float c0 = fmaxf(fabsf((a##0).x), fabsf((a##0).y));                    \
  float c1 = fmaxf(fabsf((a##1).x), fabsf((a##1).y));                   \
  float c2 = fmaxf(fabsf((a##2).x), fabsf((a##2).y));                   \
  float c3 = fmaxf(fabsf((a##3).x), fabsf((a##3).y));                   \
  float c4 = fmaxf(fabsf((a##4).x), fabsf((a##4).y));                   \
  float c5 = fmaxf(fabsf((a##5).x), fabsf((a##5).y));                   \
  float c6 = fmaxf(fabsf((a##6).x), fabsf((a##6).y));                   \
  float c7 = fmaxf(fabsf((a##7).x), fabsf((a##7).y));                   \
  float c8 = fmaxf(fabsf((a##8).x), fabsf((a##8).y));                   \
  float c9 = fmaxf(fabsf((a##9).x), fabsf((a##9).y));                   \
  float c10 = fmaxf(fabsf((a##10).x), fabsf((a##10).y));                \
  float c11 = fmaxf(fabsf((a##11).x), fabsf((a##11).y));                \
  c0 = fmaxf(c0, c1); c1 = fmaxf(c2, c3);  c2 = fmaxf(c4, c5); c3 = fmaxf(c6, c7); \
  c4 = fmaxf(c8, c9); c5 = fmaxf(c10, c11); c0 = fmaxf(c0, c1); c1 = fmaxf(c2, c3); \
  c2 = fmaxf(c4, c5); c0 = fmaxf(c0, c1); c0 = fmaxf(c0, c2);           \
  n[i] = c0;                                                            \
  float C = __fdividef(MAX_SHORT, c0);                                  \
  h[i+0*length] = make_short4((short)(C*(float)(a##0).x), (short)(C*(float)(a##0).y), \
                              (short)(C*(float)(a##1).x), (short)(C*(float)(a##1).y)); \
  h[i+1*length] = make_short4((short)(C*(float)(a##2).x), (short)(C*(float)(a##2).y), \
                              (short)(C*(float)(a##3).x), (short)(C*(float)(a##3).y)); \
  h[i+2*length] = make_short4((short)(C*(float)(a##4).x), (short)(C*(float)(a##4).y), \
                              (short)(C*(float)(a##5).x), (short)(C*(float)(a##5).y)); \
  h[i+3*length] = make_short4((short)(C*(float)(a##6).x), (short)(C*(float)(a##6).y), \
                              (short)(C*(float)(a##7).x), (short)(C*(float)(a##7).y)); \
  h[i+4*length] = make_short4((short)(C*(float)(a##8).x), (short)(C*(float)(a##8).y), \
                              (short)(C*(float)(a##9).x), (short)(C*(float)(a##9).y)); \
  h[i+5*length] = make_short4((short)(C*(float)(a##10).x), (short)(C*(float)(a##10).y), \
                              (short)(C*(float)(a##11).x), (short)(C*(float)(a##11).y));}

Definition at line 347 of file blas_quda.cu.

#define CONSTRUCT_HALF_SPINOR_FROM_DOUBLE_ST	(	h,
		n,
		a,
		length
	)

Value:

{float c0 = fmaxf(fabsf((a##0).x), fabsf((a##0).y));            \
        float c1 = fmaxf(fabsf((a##1).x), fabsf((a##1).y));             \
        float c2 = fmaxf(fabsf((a##2).x), fabsf((a##2).y));             \
        c0 = fmaxf(c0, c1); c0 = fmaxf(c0, c2);                         \
        n[i] = c0;                                                      \
        float C = __fdividef(MAX_SHORT, c0);                            \
        h[i+0*length] = make_short2((short)(C*(float)(a##0).x), (short)(C*(float)(a##0).y)); \
        h[i+1*length] = make_short2((short)(C*(float)(a##1).x), (short)(C*(float)(a##1).y)); \
        h[i+2*length] = make_short2((short)(C*(float)(a##2).x), (short)(C*(float)(a##2).y));}

Definition at line 389 of file blas_quda.cu.

#define CONSTRUCT_HALF_SPINOR_FROM_SINGLE	(	h,
		n,
		a,
		length
	)

Value:

{               \
    float c0 = fast_abs_max(a##0);                                      \
    float c1 = fast_abs_max(a##1);                                      \
    c0 = FAST_MAX(c0, c1);                                              \
    float c2 = fast_abs_max(a##2);                                      \
    float c3 = fast_abs_max(a##3);                                      \
    c1 = FAST_MAX(c2, c3);                                              \
    c0 = FAST_MAX(c0, c1);                                              \
    c2 = fast_abs_max(a##4);                                            \
    c3 = fast_abs_max(a##5);                                            \
    c1 = FAST_MAX(c2, c3);                                              \
    c0 = FAST_MAX(c0, c1);                                              \
    n[i] = c0;                                                          \
    float C = __fdividef(MAX_SHORT, c0);                                \
    h[i+0*length] = make_short4((short)(C*(float)(a##0).x), (short)(C*(float)(a##0).y), \
                                (short)(C*(float)(a##0).z), (short)(C*(float)(a##0).w)); \
    h[i+1*length] = make_short4((short)(C*(float)(a##1).x), (short)(C*(float)(a##1).y), \
                                (short)(C*(float)(a##1).z), (short)(C*(float)(a##1).w)); \
    h[i+2*length] = make_short4((short)(C*(float)(a##2).x), (short)(C*(float)(a##2).y), \
                                (short)(C*(float)(a##2).z), (short)(C*(float)(a##2).w)); \
    h[i+3*length] = make_short4((short)(C*(float)(a##3).x), (short)(C*(float)(a##3).y), \
                                (short)(C*(float)(a##3).z), (short)(C*(float)(a##3).w)); \
    h[i+4*length] = make_short4((short)(C*(float)(a##4).x), (short)(C*(float)(a##4).y), \
                                (short)(C*(float)(a##4).z), (short)(C*(float)(a##4).w)); \
    h[i+5*length] = make_short4((short)(C*(float)(a##5).x), (short)(C*(float)(a##5).y), \
                                (short)(C*(float)(a##5).z), (short)(C*(float)(a##5).w));}

Definition at line 320 of file blas_quda.cu.

#define CONSTRUCT_HALF_SPINOR_FROM_SINGLE_ST	(	h,
		n,
		a,
		length
	)

Value:

{float c0 = fmaxf(fabsf((a##0).x), fabsf((a##0).y));            \
        float c1 = fmaxf(fabsf((a##1).x), fabsf((a##1).y));             \
        float c2 = fmaxf(fabsf((a##2).x), fabsf((a##2).y));             \
        c0 = fmaxf(c0, c1); c0 = fmaxf(c0, c2);                         \
        n[i] = c0;                                                      \
        float C = __fdividef(MAX_SHORT, c0);                            \
        h[i+0*length] = make_short2((short)(C*(float)(a##0).x), (short)(C*(float)(a##0).y)); \
        h[i+1*length] = make_short2((short)(C*(float)(a##1).x), (short)(C*(float)(a##1).y)); \
        h[i+2*length] = make_short2((short)(C*(float)(a##2).x), (short)(C*(float)(a##2).y));}

Definition at line 378 of file blas_quda.cu.

#define CXPAYPBZ_FLOAT2	(	X,
		a,
		Y,
		b,
		Z
	)

Value:

{float2 z;                                                             \
  z.x = X.x + a.x*Y.x; z.x -= a.y*Y.y; z.x += b.x*Z.x; z.x -= b.y*Z.y; \
  z.y = X.y + a.y*Y.x; z.y += a.x*Y.y; z.y += b.y*Z.x; z.y += b.x*Z.y; \
  Z.x = z.x; Z.y = z.y;}

Definition at line 547 of file blas_quda.cu.

#define CXPAYPBZ_FLOAT4	(	X,
		a,
		Y,
		b,
		Z
	)

Value:

{float2 z;                                                             \
  z.x = X.x + a.x*Y.x; z.x -= a.y*Y.y; z.x += b.x*Z.x; z.x -= b.y*Z.y; \
  z.y = X.y + a.y*Y.x; z.y += a.x*Y.y; z.y += b.y*Z.x; z.y += b.x*Z.y; \
  Z.x = z.x; Z.y = z.y;                                                \
  z.x = X.z + a.x*Y.z; z.x -= a.y*Y.w; z.x += b.x*Z.z; z.x -= b.y*Z.w; \
  z.y = X.w + a.y*Y.z; z.y += a.x*Y.w; z.y += b.y*Z.z; z.y += b.x*Z.w; \
  Z.z = z.x; Z.w = z.y;}

Definition at line 538 of file blas_quda.cu.

#define EMUSYNC

Definition at line 31 of file blas_quda.cu.

#define FAST_ABS_MAX	(	a,
		b
	)	fmaxf(fabsf(a), fabsf(b));

Definition at line 311 of file blas_quda.cu.

#define FAST_MAX	(	a,
		b
	)	fmaxf(a, b);

Definition at line 312 of file blas_quda.cu.

#define IMAG_DOT_FLOAT2	(	dot,
		a,
		b
	)	float dot = a.xb.y - a.yb.x;

Definition at line 432 of file blas_quda.cu.

#define IMAG_DOT_FLOAT4	(	dot,
		a,
		b
	)	float dot = a.xb.y - a.yb.x + a.zb.w - a.wb.z;

Definition at line 422 of file blas_quda.cu.

#define MXPY_FLOAT2	(	X,
		Y
	)	Y.x -= X.x; Y.y -= X.y;

Definition at line 456 of file blas_quda.cu.

#define MXPY_FLOAT4	(	X,
		Y
	)	Y.x -= X.x; Y.y -= X.y; Y.z -= X.z; Y.w -= X.w;

Definition at line 453 of file blas_quda.cu.

#define QudaSumComplex cuComplex

Definition at line 23 of file blas_quda.cu.

#define QudaSumFloat float

Definition at line 22 of file blas_quda.cu.

#define QudaSumFloat3 float3

Definition at line 24 of file blas_quda.cu.

#define READ_DOUBLE2_TEXTURE	(	x,
		i
	)	fetch_double2(x##TexDouble2, i)

Definition at line 221 of file blas_quda.cu.

#define READ_FLOAT2_TEXTURE	(	x,
		i
	)	tex1Dfetch(x##TexSingle2, i)

Definition at line 224 of file blas_quda.cu.

#define READ_HALF_SPINOR	(	a,
		tex,
		length
	)

Value:

float4 a##0 = tex1Dfetch(tex, i + 0*length);                            \
  float4 a##1 = tex1Dfetch(tex, i + 1*length);                          \
  float4 a##2 = tex1Dfetch(tex, i + 2*length);                          \
  float4 a##3 = tex1Dfetch(tex, i + 3*length);                          \
  float4 a##4 = tex1Dfetch(tex, i + 4*length);                          \
  float4 a##5 = tex1Dfetch(tex, i + 5*length);                          \
  float a##c = a##N[i];

Definition at line 278 of file blas_quda.cu.

#define READ_HALF_SPINOR_ST	(	a,
		tex,
		length
	)

Value:

float2 a##0 = tex1Dfetch(tex, i + 0*length);                            \
  float2 a##1 = tex1Dfetch(tex, i + 1*length);                          \
  float2 a##2 = tex1Dfetch(tex, i + 2*length);                          \
  float a##c = a##N[i];

Definition at line 287 of file blas_quda.cu.

#define READ_HALF_SPINOR_TEX	(	a,
		tex,
		texNorm,
		length
	)

Value:

float a##c = tex1Dfetch(texNorm, i);                                    \
  float4 a##0 = tex1Dfetch(tex, i + 0*length);                          \
  float4 a##1 = tex1Dfetch(tex, i + 1*length);                          \
  float4 a##2 = tex1Dfetch(tex, i + 2*length);                          \
  float4 a##3 = tex1Dfetch(tex, i + 3*length);                          \
  float4 a##4 = tex1Dfetch(tex, i + 4*length);                          \
  float4 a##5 = tex1Dfetch(tex, i + 5*length);                          \

Definition at line 269 of file blas_quda.cu.

#define REAL_DOT_FLOAT2	(	dot,
		a,
		b
	)	float dot = a.xb.x + a.yb.y;

Definition at line 418 of file blas_quda.cu.

#define REAL_DOT_FLOAT4	(	dot,
		a,
		b
	)	float dot = a.xb.x + a.yb.y + a.zb.z + a.wb.w;

Definition at line 408 of file blas_quda.cu.

#define RECONSTRUCT_HALF_SPINOR	(	a,
		texHalf,
		texNorm,
		length
	)

Value:

float a##c = tex1Dfetch(texNorm, i);                                    \
  float4 a##0 = tex1Dfetch(texHalf, i + 0*length);                      \
  float4 a##1 = tex1Dfetch(texHalf, i + 1*length);                      \
  float4 a##2 = tex1Dfetch(texHalf, i + 2*length);                      \
  float4 a##3 = tex1Dfetch(texHalf, i + 3*length);                      \
  float4 a##4 = tex1Dfetch(texHalf, i + 4*length);                      \
  float4 a##5 = tex1Dfetch(texHalf, i + 5*length);                      \
  a##0 *= a##c;                                                         \
  a##1 *= a##c;                                                         \
  a##2 *= a##c;                                                         \
  a##3 *= a##c;                                                         \
  a##4 *= a##c;                                                         \
  a##5 *= a##c;

Definition at line 235 of file blas_quda.cu.

#define RECONSTRUCT_HALF_SPINOR_ST	(	a,
		texHalf,
		texNorm,
		length
	)

Value:

float a##c = tex1Dfetch(texNorm, i);                                    \
  float2 a##0 = tex1Dfetch(texHalf, i + 0*length);                      \
  float2 a##1 = tex1Dfetch(texHalf, i + 1*length);                      \
  float2 a##2 = tex1Dfetch(texHalf, i + 2*length);                      \
  (a##0) *= a##c;                                                               \
  (a##1) *= a##c;                                                               \
  (a##2) *= a##c;

Definition at line 250 of file blas_quda.cu.