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 59 /** 60 * \addtogroup UtilIterator 61 * @{ 62 */ 63 64 65 /** 66 * \brief A random-access input wrapper for dereferencing array values using a PTX cache load modifier. 67 * 68 * \par Overview 69 * - CacheModifiedInputIteratorTis a random-access input iterator that wraps a native 70 * device pointer of type <tt>ValueType*</tt>. \p ValueType references are 71 * made by reading \p ValueType values through loads modified by \p MODIFIER. 72 * - Can be used to load any data type from memory using PTX cache load modifiers (e.g., "LOAD_LDG", 73 * "LOAD_CG", "LOAD_CA", "LOAD_CS", "LOAD_CV", etc.). 74 * - Can be constructed, manipulated, and exchanged within and between host and device 75 * functions, but can only be dereferenced within device functions. 76 * - Compatible with Thrust API v1.7 or newer. 77 * 78 * \par Snippet 79 * The code snippet below illustrates the use of \p CacheModifiedInputIteratorTto 80 * dereference a device array of double using the "ldg" PTX load modifier 81 * (i.e., load values through texture cache). 82 * \par 83 * \code 84 * #include <cub/cub.cuh> // or equivalently <cub/iterator/cache_modified_input_iterator.cuh> 85 * 86 * // Declare, allocate, and initialize a device array 87 * double *d_in; // e.g., [8.0, 6.0, 7.0, 5.0, 3.0, 0.0, 9.0] 88 * 89 * // Create an iterator wrapper 90 * cub::CacheModifiedInputIterator<cub::LOAD_LDG, double> itr(d_in); 91 * 92 * // Within device code: 93 * printf("%f\n", itr[0]); // 8.0 94 * printf("%f\n", itr[1]); // 6.0 95 * printf("%f\n", itr[6]); // 9.0 96 * 97 * \endcode 98 * 99 * \tparam CacheLoadModifier The cub::CacheLoadModifier to use when accessing data 100 * \tparam ValueType The value type of this iterator 101 * \tparam OffsetT The difference type of this iterator (Default: \p ptrdiff_t) 102 */ 103 template < 104 CacheLoadModifier MODIFIER, 105 typename ValueType, 106 typename OffsetT = ptrdiff_t> 107 class CacheModifiedInputIterator 108 { 109 public: 110 111 // Required iterator traits 112 typedef CacheModifiedInputIterator self_type; ///< My own type 113 typedef OffsetT difference_type; ///< Type to express the result of subtracting one iterator from another 114 typedef ValueType value_type; ///< The type of the element the iterator can point to 115 typedef ValueType* pointer; ///< The type of a pointer to an element the iterator can point to 116 typedef ValueType reference; ///< The type of a reference to an element the iterator can point to 117 118 #if (THRUST_VERSION >= 100700) 119 // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods 120 typedef typename thrust::detail::iterator_facade_category< 121 thrust::device_system_tag, 122 thrust::random_access_traversal_tag, 123 value_type, 124 reference 125 >::type iterator_category; ///< The iterator category 126 #else 127 typedef std::random_access_iterator_tag iterator_category; ///< The iterator category 128 #endif // THRUST_VERSION 129 130 131 public: 132 133 /// Wrapped native pointer 134 ValueType* ptr; 135 136 /// Constructor 137 template <typename QualifiedValueType> CacheModifiedInputIterator(QualifiedValueType * ptr)138 __host__ __device__ __forceinline__ CacheModifiedInputIterator( 139 QualifiedValueType* ptr) ///< Native pointer to wrap 140 : 141 ptr(const_cast<typename RemoveQualifiers<QualifiedValueType>::Type *>(ptr)) 142 {} 143 144 /// Postfix increment operator ++(int)145 __host__ __device__ __forceinline__ self_type operator++(int) 146 { 147 self_type retval = *this; 148 ptr++; 149 return retval; 150 } 151 152 /// Prefix increment operator ++()153 __host__ __device__ __forceinline__ self_type operator++() 154 { 155 ptr++; 156 return *this; 157 } 158 159 /// Indirection operator *() const160 __device__ __forceinline__ reference operator*() const 161 { 162 return ThreadLoad<MODIFIER>(ptr); 163 } 164 165 /// Addition 166 template <typename Distance> operator +(Distance n) const167 __host__ __device__ __forceinline__ self_type operator+(Distance n) const 168 { 169 self_type retval(ptr + n); 170 return retval; 171 } 172 173 /// Addition assignment 174 template <typename Distance> operator +=(Distance n)175 __host__ __device__ __forceinline__ self_type& operator+=(Distance n) 176 { 177 ptr += n; 178 return *this; 179 } 180 181 /// Subtraction 182 template <typename Distance> operator -(Distance n) const183 __host__ __device__ __forceinline__ self_type operator-(Distance n) const 184 { 185 self_type retval(ptr - n); 186 return retval; 187 } 188 189 /// Subtraction assignment 190 template <typename Distance> operator -=(Distance n)191 __host__ __device__ __forceinline__ self_type& operator-=(Distance n) 192 { 193 ptr -= n; 194 return *this; 195 } 196 197 /// Distance operator -(self_type other) const198 __host__ __device__ __forceinline__ difference_type operator-(self_type other) const 199 { 200 return ptr - other.ptr; 201 } 202 203 /// Array subscript 204 template <typename Distance> operator [](Distance n) const205 __device__ __forceinline__ reference operator[](Distance n) const 206 { 207 return ThreadLoad<MODIFIER>(ptr + n); 208 } 209 210 /// Structure dereference operator ->()211 __device__ __forceinline__ pointer operator->() 212 { 213 return &ThreadLoad<MODIFIER>(ptr); 214 } 215 216 /// Equal to operator ==(const self_type & rhs)217 __host__ __device__ __forceinline__ bool operator==(const self_type& rhs) 218 { 219 return (ptr == rhs.ptr); 220 } 221 222 /// Not equal to operator !=(const self_type & rhs)223 __host__ __device__ __forceinline__ bool operator!=(const self_type& rhs) 224 { 225 return (ptr != rhs.ptr); 226 } 227 228 /// ostream operator operator <<(std::ostream & os,const self_type &)229 friend std::ostream& operator<<(std::ostream& os, const self_type& /*itr*/) 230 { 231 return os; 232 } 233 }; 234 235 236 237 /** @} */ // end group UtilIterator 238 239 } // CUB namespace 240 CUB_NS_POSTFIX // Optional outer namespace(s) 241