cub::TexObjInputIterator#

template<typename T, typename OffsetT = ptrdiff_t>
class TexObjInputIterator#

A random-access input wrapper for dereferencing array values through texture cache.

Uses newer Kepler-style texture objects.

Overview

  • TexObjInputIterator wraps 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, TexObjInputIterator iterators 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 TexObjInputIterator to 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

using self_type = TexObjInputIterator#

My own type.

using difference_type = OffsetT#

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

using value_type = T#

The type of the element the iterator can point to.

using pointer = T*#

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

using reference = T#

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

using iterator_category = typename THRUST_NS_QUALIFIER::detail::iterator_facade_category<THRUST_NS_QUALIFIER::device_system_tag, THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference>::type#

The iterator category.

Public Functions

inline TexObjInputIterator()#

Constructor.

template<typename QualifiedT>
inline cudaError_t BindTexture(
QualifiedT *ptr,
size_t bytes,
size_t tex_offset = 0,
)#

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

inline cudaError_t UnbindTexture()#

Unbind this iterator from its texture reference.

inline self_type operator++(int)#

Postfix increment.

inline self_type operator++()#

Prefix increment.

inline reference operator*() const#

Indirection.

template<typename Distance>
inline self_type operator+(
Distance n,
) const#

Addition.

template<typename Distance>
inline self_type &operator+=(Distance n)#

Addition assignment.

template<typename Distance>
inline self_type operator-(
Distance n,
) const#

Subtraction.

template<typename Distance>
inline self_type &operator-=(Distance n)#

Subtraction assignment.

inline difference_type operator-(self_type other) const#

Distance.

template<typename Distance>
inline reference operator[](
Distance n,
) const#

Array subscript.

inline pointer operator->()#

Structure dereference.

inline bool operator==(const self_type &rhs) const#

Equal to.

inline bool operator!=(const self_type &rhs) const#

Not equal to.

inline friend ::std::ostream &operator<<(
::std::ostream &os,
const self_type &itr,
)#

ostream operator