|
/****************************************************************************** |
|
* Copyright (c) 2011, Duane Merrill. All rights reserved. |
|
* Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved. |
|
* |
|
* Redistribution and use in source and binary forms, with or without |
|
* modification, are permitted provided that the following conditions are met: |
|
* * Redistributions of source code must retain the above copyright |
|
* notice, this list of conditions and the following disclaimer. |
|
* * Redistributions in binary form must reproduce the above copyright |
|
* notice, this list of conditions and the following disclaimer in the |
|
* documentation and/or other materials provided with the distribution. |
|
* * Neither the name of the NVIDIA CORPORATION nor the |
|
* names of its contributors may be used to endorse or promote products |
|
* derived from this software without specific prior written permission. |
|
* |
|
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND |
|
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED |
|
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE |
|
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY |
|
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES |
|
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; |
|
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND |
|
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT |
|
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS |
|
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
|
* |
|
******************************************************************************/ |
|
|
|
/** |
|
* \file |
|
* Random-access iterator types |
|
*/ |
|
|
|
#pragma once |
|
|
|
#include <iterator> |
|
#include <iostream> |
|
|
|
#include "../thread/thread_load.cuh" |
|
#include "../thread/thread_store.cuh" |
|
#include "../config.cuh" |
|
#include "../util_device.cuh" |
|
|
|
#if (THRUST_VERSION >= 100700) |
|
// This iterator is compatible with Thrust API 1.7 and newer |
|
#include <thrust/iterator/iterator_facade.h> |
|
#include <thrust/iterator/iterator_traits.h> |
|
#endif // THRUST_VERSION |
|
|
|
|
|
CUB_NAMESPACE_BEGIN |
|
|
|
/** |
|
* \addtogroup UtilIterator |
|
* @{ |
|
*/ |
|
|
|
|
|
/** |
|
* \brief A random-access input wrapper for transforming dereferenced values. |
|
* |
|
* \par Overview |
|
* - TransformInputIteratorTwraps a unary conversion functor of type \p |
|
* ConversionOp and a random-access input iterator of type <tt>InputIteratorT</tt>, |
|
* using the former to produce references of type \p ValueType from the latter. |
|
* - Can be used with any data type. |
|
* - Can be constructed, manipulated, and exchanged within and between host and device |
|
* functions. Wrapped host memory can only be dereferenced on the host, and wrapped |
|
* device memory can only be dereferenced on the device. |
|
* - Compatible with Thrust API v1.7 or newer. |
|
* |
|
* \par Snippet |
|
* The code snippet below illustrates the use of \p TransformInputIteratorTto |
|
* dereference an array of integers, tripling the values and converting them to doubles. |
|
* \par |
|
* \code |
|
* #include <cub/cub.cuh> // or equivalently <cub/iterator/transform_input_iterator.cuh> |
|
* |
|
* // Functor for tripling integer values and converting to doubles |
|
* struct TripleDoubler |
|
* { |
|
* __host__ __device__ __forceinline__ |
|
* double operator()(const int &a) const { |
|
* return double(a * 3); |
|
* } |
|
* }; |
|
* |
|
* // Declare, allocate, and initialize a device array |
|
* int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9] |
|
* TripleDoubler conversion_op; |
|
* |
|
* // Create an iterator wrapper |
|
* cub::TransformInputIterator<double, TripleDoubler, int*> itr(d_in, conversion_op); |
|
* |
|
* // Within device code: |
|
* printf("%f\n", itr[0]); // 24.0 |
|
* printf("%f\n", itr[1]); // 18.0 |
|
* printf("%f\n", itr[6]); // 27.0 |
|
* |
|
* \endcode |
|
* |
|
* \tparam ValueType The value type of this iterator |
|
* \tparam ConversionOp Unary functor type for mapping objects of type \p InputType to type \p ValueType. Must have member <tt>ValueType operator()(const InputType &datum)</tt>. |
|
* \tparam InputIteratorT The type of the wrapped input iterator |
|
* \tparam OffsetT The difference type of this iterator (Default: \p ptrdiff_t) |
|
* |
|
*/ |
|
template < |
|
typename ValueType, |
|
typename ConversionOp, |
|
typename InputIteratorT, |
|
typename OffsetT = ptrdiff_t> |
|
class TransformInputIterator |
|
{ |
|
public: |
|
|
|
// Required iterator traits |
|
typedef TransformInputIterator self_type; ///< My own type |
|
typedef OffsetT difference_type; ///< Type to express the result of subtracting one iterator from another |
|
typedef ValueType value_type; ///< The type of the element the iterator can point to |
|
typedef ValueType* pointer; ///< The type of a pointer to an element the iterator can point to |
|
typedef ValueType reference; ///< The type of a reference to an element the iterator can point to |
|
|
|
#if (THRUST_VERSION >= 100700) |
|
// Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods |
|
typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< |
|
THRUST_NS_QUALIFIER::any_system_tag, |
|
THRUST_NS_QUALIFIER::random_access_traversal_tag, |
|
value_type, |
|
reference |
|
>::type iterator_category; ///< The iterator category |
|
#else |
|
typedef std::random_access_iterator_tag iterator_category; ///< The iterator category |
|
#endif // THRUST_VERSION |
|
|
|
private: |
|
|
|
ConversionOp conversion_op; |
|
InputIteratorT input_itr; |
|
|
|
public: |
|
|
|
/// Constructor |
|
__host__ __device__ __forceinline__ TransformInputIterator( |
|
InputIteratorT input_itr, ///< Input iterator to wrap |
|
ConversionOp conversion_op) ///< Conversion functor to wrap |
|
: |
|
conversion_op(conversion_op), |
|
input_itr(input_itr) |
|
{} |
|
|
|
/// Postfix increment |
|
__host__ __device__ __forceinline__ self_type operator++(int) |
|
{ |
|
self_type retval = *this; |
|
input_itr++; |
|
return retval; |
|
} |
|
|
|
/// Prefix increment |
|
__host__ __device__ __forceinline__ self_type operator++() |
|
{ |
|
input_itr++; |
|
return *this; |
|
} |
|
|
|
/// Indirection |
|
__host__ __device__ __forceinline__ reference operator*() const |
|
{ |
|
return conversion_op(*input_itr); |
|
} |
|
|
|
/// Addition |
|
template <typename Distance> |
|
__host__ __device__ __forceinline__ self_type operator+(Distance n) const |
|
{ |
|
self_type retval(input_itr + n, conversion_op); |
|
return retval; |
|
} |
|
|
|
/// Addition assignment |
|
template <typename Distance> |
|
__host__ __device__ __forceinline__ self_type& operator+=(Distance n) |
|
{ |
|
input_itr += n; |
|
return *this; |
|
} |
|
|
|
/// Subtraction |
|
template <typename Distance> |
|
__host__ __device__ __forceinline__ self_type operator-(Distance n) const |
|
{ |
|
self_type retval(input_itr - n, conversion_op); |
|
return retval; |
|
} |
|
|
|
/// Subtraction assignment |
|
template <typename Distance> |
|
__host__ __device__ __forceinline__ self_type& operator-=(Distance n) |
|
{ |
|
input_itr -= n; |
|
return *this; |
|
} |
|
|
|
/// Distance |
|
__host__ __device__ __forceinline__ difference_type operator-(self_type other) const |
|
{ |
|
return input_itr - other.input_itr; |
|
} |
|
|
|
/// Array subscript |
|
template <typename Distance> |
|
__host__ __device__ __forceinline__ reference operator[](Distance n) const |
|
{ |
|
return conversion_op(input_itr[n]); |
|
} |
|
|
|
/// Equal to |
|
__host__ __device__ __forceinline__ bool operator==(const self_type& rhs) |
|
{ |
|
return (input_itr == rhs.input_itr); |
|
} |
|
|
|
/// Not equal to |
|
__host__ __device__ __forceinline__ bool operator!=(const self_type& rhs) |
|
{ |
|
return (input_itr != rhs.input_itr); |
|
} |
|
|
|
/// ostream operator |
|
friend std::ostream& operator<<(std::ostream& os, const self_type& /* itr */) |
|
{ |
|
return os; |
|
} |
|
}; |
|
|
|
|
|
|
|
/** @} */ // end group UtilIterator |
|
|
|
CUB_NAMESPACE_END |
|
|