A random-access input wrapper for dereferencing array values through texture cache. Uses newer Kepler-style texture objects. More...

Detailed Description

A random-access input wrapper for dereferencing array values through texture cache. Uses newer Kepler-style texture objects.

Overview

TexObjInputIteratorTwraps a native device pointer of type ValueType*. References to elements are to be loaded through texture cache.
Can be used to load any data type from memory through texture cache.
Can be manipulated and exchanged within and between host and device functions, can only be constructed within host functions, and can only be dereferenced within device functions.
With regard to nested/dynamic parallelism, TexObjInputIteratorTiterators may only be created by the host thread, but can be used by any descendant kernel.
Compatible with Thrust API v1.7 or newer.

Snippet: The code snippet below illustrates the use of TexRefInputIteratorTto dereference a device array of doubles through texture cache.

: #include <cub/cub.cuh> // or equivalently <cub/iterator/tex_obj_input_iterator.cuh>
// Declare, allocate, and initialize a device array
int num_items; // e.g., 7
double *d_in; // e.g., [8.0, 6.0, 7.0, 5.0, 3.0, 0.0, 9.0]
// Create an iterator wrapper
cub::TexObjInputIterator<double> itr;
itr.BindTexture(d_in, sizeof(double) * num_items);
...
// Within device code:
printf("%f\n", itr[0]); // 8.0
printf("%f\n", itr[1]); // 6.0
printf("%f\n", itr[6]); // 9.0
...
itr.UnbindTexture();

Template Parameters

T	The value type of this iterator
OffsetT	The difference type of this iterator (Default: `ptrdiff_t`)

Public Types
typedef TexObjInputIterator	self_type
	My own type.

typedef OffsetT	difference_type
	Type to express the result of subtracting one iterator from another.

typedef T	value_type
	The type of the element the iterator can point to.

typedef T *	pointer
	The type of a pointer to an element the iterator can point to.

typedef T	reference
	The type of a reference to an element the iterator can point to.

typedef std::random_access_iterator_tag	iterator_category
	The iterator category.

Public Member Functions
__host__ __device__ __forceinline__	TexObjInputIterator ()
	Constructor.

template<typename QualifiedT >
cudaError_t	BindTexture (QualifiedT *ptr, size_t bytes=size_t(-1), size_t tex_offset=0)
	Use this iterator to bind `ptr` with a texture reference. More...

cudaError_t	UnbindTexture ()
	Unbind this iterator from its texture reference.

__host__ __device__ __forceinline__ self_type	operator++ (int)
	Postfix increment.

__host__ __device__ __forceinline__ self_type	operator++ ()
	Prefix increment.

__host__ __device__ __forceinline__ reference	operator * () const
	Indirection.

template<typename Distance >
__host__ __device__ __forceinline__ self_type	operator+ (Distance n) const
	Addition.

template<typename Distance >
__host__ __device__ __forceinline__ self_type &	operator+= (Distance n)
	Addition assignment.

template<typename Distance >
__host__ __device__ __forceinline__ self_type	operator- (Distance n) const
	Subtraction.

template<typename Distance >
__host__ __device__ __forceinline__ self_type &	operator-= (Distance n)
	Subtraction assignment.

__host__ __device__ __forceinline__ difference_type	operator- (self_type other) const
	Distance.

template<typename Distance >
__host__ __device__ __forceinline__ reference	operator[] (Distance n) const
	Array subscript.

__host__ __device__ __forceinline__ pointer	operator-> ()
	Structure dereference.

__host__ __device__ __forceinline__ bool	operator== (const self_type &rhs)
	Equal to.

__host__ __device__ __forceinline__ bool	operator!= (const self_type &rhs)
	Not equal to.

Private Types
enum	{ TEXTURE_MULTIPLE = sizeof(T) / sizeof(TextureWord) }

typedef UnitWord< T >::TextureWord	TextureWord

Friends
std::ostream &	operator<< (std::ostream &os, const self_type &itr)
	ostream operator

Member Function Documentation

template<typename T, typename OffsetT = ptrdiff_t>

template<typename QualifiedT >

cudaError_t cub::TexObjInputIterator< T, OffsetT >::BindTexture	(	QualifiedT *	ptr,
		size_t	bytes = `size_t(-1)`,
		size_t	tex_offset = `0`
	)

inline

Use this iterator to bind ptr with a texture reference.

Parameters

ptr	Native pointer to wrap that is aligned to cudaDeviceProp::textureAlignment
bytes	Number of bytes in the range
tex_offset	OffsetT (in items) from `ptr` denoting the position of the iterator

The documentation for this class was generated from the following file: