texture.h File Reference

#include <convert.h>

Include dependency graph for texture.h:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes
class	Texture< OutputType, InputType, tex_id >
class	SpinorTexture< RegType, StoreType, N, tex_id >
class	Spinor< RegType, 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	SPINOR   ST::spinor

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

DECL_TEX

#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 174 of file texture.h.

DEF_ALL

#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 228 of file texture.h.

DEF_BIND_UNBIND

#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 183 of file texture.h.

DEF_BIND_UNBIND_FETCH

#define DEF_BIND_UNBIND_FETCH	(		outtype,
			intype,
			id
	)

Value:

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

Definition at line 223 of file texture.h.

DEF_FETCH

#define DEF_FETCH   DEF_FETCH_TEX

Definition at line 202 of file texture.h.

DEF_FETCH_DBLE

#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 209 of file texture.h.

DEF_FETCH_DBLE_MIXED

#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 217 of file texture.h.

DEF_FETCH_DIRECT

#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 194 of file texture.h.

DEF_FETCH_TEX

#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 189 of file texture.h.

MAX_TEX_ID

#define MAX_TEX_ID   4

Definition at line 97 of file texture.h.

Referenced by SpinorTexture< RegType, StoreType, N, tex_id >::load(), SpinorTexture< RegType, StoreType, N, tex_id >::loadGhost(), and Texture< InterType, StoreType, -1 >::~Texture().

MAX_TEXELS

#define MAX_TEXELS   (1<<27)

Definition at line 95 of file texture.h.

SPINOR

#define SPINOR   ST::spinor

Definition at line 539 of file texture.h.

Referenced by Spinor< RegType, StoreType, N, write, tex_id >::backup(), Spinor< RegType, StoreType, N, write, tex_id >::restore(), Spinor< RegType, StoreType, N, write, tex_id >::save(), and Spinor< RegType, StoreType, N, write, tex_id >::V().

Function Documentation

checkTypes()

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 279 of file texture.h.

References errorQuda.

fetch_double()

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

Definition at line 86 of file texture.h.

fetch_double2()

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

Definition at line 88 of file texture.h.

store_norm()

template<typename FloatN , int M>

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

inline

Definition at line 300 of file texture.h.

References c, fmaxf(), fused_exterior_ndeg_tm_dslash_cuda_gen::i, quda::max_fabs(), MAX_SHORT, quda::norm(), and x.

Here is the call graph for this function:

Variable Documentation

tex_id_table

bool tex_id_table[MAX_TEX_ID] = { }

Definition at line 100 of file texture.h.

Referenced by Texture< InterType, StoreType, -1 >::~Texture().