1 /****************************************************************************** 2 * Copyright (c) 2011, Duane Merrill. All rights reserved. 3 * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved. 4 * 5 * Redistribution and use in source and binary forms, with or without 6 * modification, are permitted provided that the following conditions are met: 7 * * Redistributions of source code must retain the above copyright 8 * notice, this list of conditions and the following disclaimer. 9 * * Redistributions in binary form must reproduce the above copyright 10 * notice, this list of conditions and the following disclaimer in the 11 * documentation and/or other materials provided with the distribution. 12 * * Neither the name of the NVIDIA CORPORATION nor the 13 * names of its contributors may be used to endorse or promote products 14 * derived from this software without specific prior written permission. 15 * 16 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND 17 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED 18 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE 19 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY 20 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES 21 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; 22 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND 23 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT 24 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS 25 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. 26 * 27 ******************************************************************************/ 28 29 /** 30 * \file 31 * Random-access iterator types 32 */ 33 34 #pragma once 35 36 #include <iterator> 37 #include <iostream> 38 39 #include "../thread/thread_load.cuh" 40 #include "../thread/thread_store.cuh" 41 #include "../util_device.cuh" 42 #include "../util_namespace.cuh" 43 44 #if (THRUST_VERSION >= 100700) 45 // This iterator is compatible with Thrust API 1.7 and newer 46 #include <thrust/iterator/iterator_facade.h> 47 #include <thrust/iterator/iterator_traits.h> 48 #endif // THRUST_VERSION 49 50 51 /// Optional outer namespace(s) 52 CUB_NS_PREFIX 53 54 /// CUB namespace 55 namespace cub { 56 57 /** 58 * \addtogroup UtilIterator 59 * @{ 60 */ 61 62 /** 63 * \brief A random-access input generator for dereferencing a sequence of incrementing integer values. 64 * 65 * \par Overview 66 * - After initializing a CountingInputIteratorTto a certain integer \p base, read references 67 * at \p offset will return the value \p base + \p offset. 68 * - Can be constructed, manipulated, dereferenced, and exchanged within and between host and device 69 * functions. 70 * - Compatible with Thrust API v1.7 or newer. 71 * 72 * \par Snippet 73 * The code snippet below illustrates the use of \p CountingInputIteratorTto 74 * dereference a sequence of incrementing integers. 75 * \par 76 * \code 77 * #include <cub/cub.cuh> // or equivalently <cub/iterator/counting_input_iterator.cuh> 78 * 79 * cub::CountingInputIterator<int> itr(5); 80 * 81 * printf("%d\n", itr[0]); // 5 82 * printf("%d\n", itr[1]); // 6 83 * printf("%d\n", itr[2]); // 7 84 * printf("%d\n", itr[50]); // 55 85 * 86 * \endcode 87 * 88 * \tparam ValueType The value type of this iterator 89 * \tparam OffsetT The difference type of this iterator (Default: \p ptrdiff_t) 90 */ 91 template < 92 typename ValueType, 93 typename OffsetT = ptrdiff_t> 94 class CountingInputIterator 95 { 96 public: 97 98 // Required iterator traits 99 typedef CountingInputIterator self_type; ///< My own type 100 typedef OffsetT difference_type; ///< Type to express the result of subtracting one iterator from another 101 typedef ValueType value_type; ///< The type of the element the iterator can point to 102 typedef ValueType* pointer; ///< The type of a pointer to an element the iterator can point to 103 typedef ValueType reference; ///< The type of a reference to an element the iterator can point to 104 105 #if (THRUST_VERSION >= 100700) 106 // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods 107 typedef typename thrust::detail::iterator_facade_category< 108 thrust::any_system_tag, 109 thrust::random_access_traversal_tag, 110 value_type, 111 reference 112 >::type iterator_category; ///< The iterator category 113 #else 114 typedef std::random_access_iterator_tag iterator_category; ///< The iterator category 115 #endif // THRUST_VERSION 116 117 private: 118 119 ValueType val; 120 121 public: 122 123 /// Constructor CountingInputIterator(const ValueType & val)124 __host__ __device__ __forceinline__ CountingInputIterator( 125 const ValueType &val) ///< Starting value for the iterator instance to report 126 : 127 val(val) 128 {} 129 130 /// Postfix increment operator ++(int)131 __host__ __device__ __forceinline__ self_type operator++(int) 132 { 133 self_type retval = *this; 134 val++; 135 return retval; 136 } 137 138 /// Prefix increment operator ++()139 __host__ __device__ __forceinline__ self_type operator++() 140 { 141 val++; 142 return *this; 143 } 144 145 /// Indirection operator *() const146 __host__ __device__ __forceinline__ reference operator*() const 147 { 148 return val; 149 } 150 151 /// Addition 152 template <typename Distance> operator +(Distance n) const153 __host__ __device__ __forceinline__ self_type operator+(Distance n) const 154 { 155 self_type retval(val + (ValueType) n); 156 return retval; 157 } 158 159 /// Addition assignment 160 template <typename Distance> operator +=(Distance n)161 __host__ __device__ __forceinline__ self_type& operator+=(Distance n) 162 { 163 val += (ValueType) n; 164 return *this; 165 } 166 167 /// Subtraction 168 template <typename Distance> operator -(Distance n) const169 __host__ __device__ __forceinline__ self_type operator-(Distance n) const 170 { 171 self_type retval(val - (ValueType) n); 172 return retval; 173 } 174 175 /// Subtraction assignment 176 template <typename Distance> operator -=(Distance n)177 __host__ __device__ __forceinline__ self_type& operator-=(Distance n) 178 { 179 val -= n; 180 return *this; 181 } 182 183 /// Distance operator -(self_type other) const184 __host__ __device__ __forceinline__ difference_type operator-(self_type other) const 185 { 186 return (difference_type) (val - other.val); 187 } 188 189 /// Array subscript 190 template <typename Distance> operator [](Distance n) const191 __host__ __device__ __forceinline__ reference operator[](Distance n) const 192 { 193 return val + (ValueType) n; 194 } 195 196 /// Structure dereference operator ->()197 __host__ __device__ __forceinline__ pointer operator->() 198 { 199 return &val; 200 } 201 202 /// Equal to operator ==(const self_type & rhs)203 __host__ __device__ __forceinline__ bool operator==(const self_type& rhs) 204 { 205 return (val == rhs.val); 206 } 207 208 /// Not equal to operator !=(const self_type & rhs)209 __host__ __device__ __forceinline__ bool operator!=(const self_type& rhs) 210 { 211 return (val != rhs.val); 212 } 213 214 /// ostream operator operator <<(std::ostream & os,const self_type & itr)215 friend std::ostream& operator<<(std::ostream& os, const self_type& itr) 216 { 217 os << "[" << itr.val << "]"; 218 return os; 219 } 220 221 }; 222 223 224 225 /** @} */ // end group UtilIterator 226 227 } // CUB namespace 228 CUB_NS_POSTFIX // Optional outer namespace(s) 229