texture.h File Reference

#include <quda_internal.h>
#include <color_spinor_field.h>
#include <convert.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	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_BIND_UNBIND_FETCH(outtype, intype, id)
#define	DEF_ALL(id)
#define	MAX_TEX_ID   4
#define	REG_LENGTH   (sizeof(RegType) / sizeof(((RegType*)0)->x))
#define	IS_SHORT(type)   (sizeof( ((type*)0)->x ) == sizeof(short) )

Functions
__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)

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 136 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)

Definition at line 182 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 145 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 177 of file texture.h.

#define DEF_FETCH   DEF_FETCH_TEX

Definition at line 164 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 171 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 156 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 151 of file texture.h.

#define IS_SHORT

(

type

)

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

Definition at line 264 of file texture.h.

#define MAX_TEX_ID   4

Definition at line 202 of file texture.h.

#define MAX_TEXELS   (1<<27)

Definition at line 77 of file texture.h.

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

Definition at line 261 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 229 of file texture.h.

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

Definition at line 70 of file texture.h.

template<typename FloatN , int M>

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

inline

Definition at line 250 of file texture.h.