cub::TexObjInputIterator

Defined in cub/iterator/tex_obj_input_iterator.cuh

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.

Friends

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

ostream operator