cub::CacheModifiedInputIterator#

template<CacheLoadModifier MODIFIER, typename ValueType, typename OffsetT = ptrdiff_t>
class CacheModifiedInputIterator#

A random-access input wrapper for dereferencing array values using a PTX cache load modifier.

Overview

  • CacheModifiedInputIterator is a random-access input iterator that wraps a native device pointer of type ValueType*. ValueType references are made by reading ValueType values through loads modified by MODIFIER.

  • Can be used to load any data type from memory using PTX cache load modifiers (e.g., “LOAD_LDG”, “LOAD_CG”, “LOAD_CA”, “LOAD_CS”, “LOAD_CV”, etc.).

  • Can be constructed, manipulated, and exchanged within and between host and device functions, but can only be dereferenced within device functions.

  • Compatible with Thrust API v1.7 or newer.

Snippet

The code snippet below illustrates the use of CacheModifiedInputIterator to dereference a device array of double using the “ldg” PTX load modifier (i.e., load values through texture cache). 

#include <cub/cub.cuh>   // or equivalently <cub/iterator/cache_modified_input_iterator.cuh>

// Declare, allocate, and initialize a device array
double *d_in;            // e.g., [8.0, 6.0, 7.0, 5.0, 3.0, 0.0, 9.0]

// Create an iterator wrapper
cub::CacheModifiedInputIterator<cub::LOAD_LDG, double> itr(d_in);

// Within device code:
printf("%f\n", itr[0]);  // 8.0
printf("%f\n", itr[1]);  // 6.0
printf("%f\n", itr[6]);  // 9.0

Template Parameters:

  • CacheLoadModifier – The cub::CacheLoadModifier to use when accessing data

  • ValueType – The value type of this iterator

  • OffsetT – The difference type of this iterator (Default: ptrdiff_t)

Public Types

using self_type = CacheModifiedInputIterator#

My own type.

using difference_type = OffsetT#

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

using value_type = ValueType#

The type of the element the iterator can point to.

using pointer = ValueType*#

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

using reference = ValueType#

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#

Public Functions

template<typename QualifiedValueType>
inline CacheModifiedInputIterator(
QualifiedValueType *ptr,
)#

Constructor.

Parameters:

ptr – Native pointer to wrap

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&,
)#

ostream operator

Public Members

ValueType *ptr#

Wrapped native pointer.