texture.h File Reference

#include <quda_internal.h>
#include <color_spinor_field.h>
#include <convert.h>
#include <register_traits.h>
#include <float_vector.h>

Go to the source code of this file.

Classes
class	Texture< OutputType, InputType, tex_id >
class	Spinor< RegType, InterType, StoreType, N, write, tex_id >

Macros
#define	MAX_TEXELS   (1<<27)
#define	MAX_TEX_ID   4
#define	DECL_TEX(id)
#define	DEF_BIND_UNBIND(outtype, intype, id)
#define	DEF_FETCH_TEX(outtype, intype, id)
#define	DEF_FETCH_DIRECT(outtype, intype, id)
#define	DEF_FETCH   DEF_FETCH_TEX
#define	DEF_FETCH_DBLE(outtype, intype, id)
#define	DEF_FETCH_DBLE_MIXED(outtype, intype, id)
#define	DEF_BIND_UNBIND_FETCH(outtype, intype, id)
#define	DEF_ALL(id)
#define	REG_LENGTH   (sizeof(RegType) / sizeof(((RegType*)0)->x))
#define	IS_SHORT(type)   (sizeof( ((type*)0)->x ) == sizeof(short) )

Functions
__inline__ __device__ double	fetch_double (texture< int2, 1 > t, int i)
__inline__ __device__ double2	fetch_double2 (texture< int4, 1 > t, int i)
template<typename RegType , typename InterType , typename StoreType >
void	checkTypes ()
template<typename FloatN , int M>
__device__ float	store_norm (float *norm, FloatN x[M], int i)

Variables
bool	tex_id_table [MAX_TEX_ID] = { }

Macro Definition Documentation

#define DECL_TEX

(

id

)

Value:

texture<short2,1,cudaReadModeNormalizedFloat> tex_short2_##id;  \
  texture<short4,1,cudaReadModeNormalizedFloat> tex_short4_##id;        \
  texture<float,1> tex_float_##id;                                      \
  texture<float2,1> tex_float2_##id;                                    \
  texture<float4,1> tex_float4_##id;                                    \
  texture<int4,1> tex_double2_##id;

Definition at line 171 of file texture.h.

#define DEF_ALL

(

id

)

Value:

DECL_TEX(id)                                    \
  DEF_BIND_UNBIND_FETCH(float2, short2, id)     \
  DEF_BIND_UNBIND_FETCH(float4, short4, id)     \
  DEF_BIND_UNBIND_FETCH(float, float, id)       \
  DEF_BIND_UNBIND_FETCH(float2, float2, id)     \
  DEF_BIND_UNBIND_FETCH(float4, float4, id)     \
  DEF_BIND_UNBIND(double2, double2, id)         \
  DEF_BIND_UNBIND(float2, double2, id)          \
  DEF_FETCH_DBLE(double2, double2, id)          \
  DEF_FETCH_DBLE(float2, double2, id)           \
  DEF_BIND_UNBIND(double2, float2, id)          \
  DEF_BIND_UNBIND(double4, float4, id)          \
  DEF_BIND_UNBIND(double2, short2, id)          \
  DEF_BIND_UNBIND(double4, short4, id)          \
  DEF_FETCH_DBLE_MIXED(double2, float2, id)     \
  DEF_FETCH_DBLE_MIXED(double4, float4, id)     \
  DEF_FETCH_DBLE_MIXED(double2, short2, id)     \
  DEF_FETCH_DBLE_MIXED(double4, short4, id)

Definition at line 225 of file texture.h.

#define DEF_BIND_UNBIND	(		outtype,
			intype,
			id
	)

Value:

template<> inline void Texture<outtype,intype,id>::bind(const intype *ptr, size_t bytes) \
    { cudaBindTexture(0,tex_##intype##_##id, ptr, bytes); }             \
  template<> inline void Texture<outtype,intype,id>::unbind() { cudaUnbindTexture(tex_##intype##_##id); }

Definition at line 180 of file texture.h.

#define DEF_BIND_UNBIND_FETCH	(		outtype,
			intype,
			id
	)

Value:

DEF_BIND_UNBIND(outtype, intype, id)                    \
  DEF_FETCH(outtype, intype, id)

Definition at line 220 of file texture.h.

#define DEF_FETCH   DEF_FETCH_TEX

Definition at line 199 of file texture.h.

#define DEF_FETCH_DBLE	(		outtype,
			intype,
			id
	)

Value:

template<> __device__ inline outtype Texture<outtype,double2,id>::fetch(unsigned int idx) \
    { outtype out; copyFloatN(out, fetch_double2(tex_double2_##id,idx)); return out; }

Definition at line 206 of file texture.h.

#define DEF_FETCH_DBLE_MIXED	(		outtype,
			intype,
			id
	)

Value:

template<> __device__ inline outtype Texture<outtype,intype,id>::fetch(unsigned int idx) \
  { outtype out; copyFloatN(out, tex1Dfetch(tex_##intype##_##id,idx)); return out; }

Definition at line 214 of file texture.h.

#define DEF_FETCH_DIRECT	(		outtype,
			intype,
			id
	)

Value:

template<> __device__ inline outtype Texture<outtype,intype,id>::fetch(unsigned int idx) \
    { outtype out; copyFloatN(out, spinor[idx]); return out; }

Definition at line 191 of file texture.h.

#define DEF_FETCH_TEX	(		outtype,
			intype,
			id
	)

Value:

template<> __device__ inline outtype Texture<outtype,intype,id>::fetch(unsigned int idx) \
    { return tex1Dfetch(tex_##intype##_##id,idx); }

Definition at line 186 of file texture.h.

#define IS_SHORT

(

type

)

(sizeof( ((type*)0)->x ) == sizeof(short) )

Definition at line 311 of file texture.h.

#define MAX_TEX_ID   4

Definition at line 99 of file texture.h.

#define MAX_TEXELS   (1<<27)

Definition at line 97 of file texture.h.

#define REG_LENGTH   (sizeof(RegType) / sizeof(((RegType*)0)->x))

Definition at line 308 of file texture.h.

Function Documentation

template<typename RegType , typename InterType , typename StoreType >

void checkTypes

(

)

Checks that the types are set correctly. The precision used in the RegType must match that of the InterType, and the ordering of the InterType must match that of the StoreType. The only exception is when half precision is used, in which case, RegType can be a double and InterType can be single (with StoreType short).

Parameters

RegType	Register type used in kernel
InterType	Intermediate format - RegType precision with StoreType ordering
StoreType	Type used to store field in memory

Definition at line 276 of file texture.h.

__inline__ __device__ double fetch_double	(	texture< int2, 1 >	t,
		int	i
	)

Definition at line 88 of file texture.h.

__inline__ __device__ double2 fetch_double2	(	texture< int4, 1 >	t,
		int	i
	)

Definition at line 90 of file texture.h.

template<typename FloatN , int M>

__device__ float store_norm ( float *  norm, FloatN  x[M], int  i  )

inline

Definition at line 297 of file texture.h.

Variable Documentation

bool tex_id_table[MAX_TEX_ID] = { }

Definition at line 102 of file texture.h.