cub::CacheModifiedInputIterator

Defined in /home/runner/work/cccl/cccl/cub/cub/iterator/cache_modified_input_iterator.cuh

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 = std::random_access_iterator_tag

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.

Public Members

ValueType *ptr

Wrapped native pointer.

Friends

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

ostream operator