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 "../util_namespace.cuh" 40 #include "../util_macro.cuh" 41 42 #if (THRUST_VERSION >= 100700) 43 // This iterator is compatible with Thrust API 1.7 and newer 44 #include <thrust/iterator/iterator_facade.h> 45 #include <thrust/iterator/iterator_traits.h> 46 #endif // THRUST_VERSION 47 48 49 /// Optional outer namespace(s) 50 CUB_NS_PREFIX 51 52 /// CUB namespace 53 namespace cub { 54 55 56 /** 57 * \addtogroup UtilIterator 58 * @{ 59 */ 60 61 62 /** 63 * \brief A discard iterator 64 */ 65 template <typename OffsetT = ptrdiff_t> 66 class DiscardOutputIterator 67 { 68 public: 69 70 // Required iterator traits 71 typedef DiscardOutputIterator self_type; ///< My own type 72 typedef OffsetT difference_type; ///< Type to express the result of subtracting one iterator from another 73 typedef void value_type; ///< The type of the element the iterator can point to 74 typedef void pointer; ///< The type of a pointer to an element the iterator can point to 75 typedef void reference; ///< The type of a reference to an element the iterator can point to 76 77 #if (THRUST_VERSION >= 100700) 78 // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods 79 typedef typename thrust::detail::iterator_facade_category< 80 thrust::any_system_tag, 81 thrust::random_access_traversal_tag, 82 value_type, 83 reference 84 >::type iterator_category; ///< The iterator category 85 #else 86 typedef std::random_access_iterator_tag iterator_category; ///< The iterator category 87 #endif // THRUST_VERSION 88 89 private: 90 91 OffsetT offset; 92 93 #if defined(_WIN32) || !defined(_WIN64) 94 // Workaround for win32 parameter-passing bug (ulonglong2 argmin DeviceReduce) 95 OffsetT pad[CUB_MAX(1, (16 / sizeof(OffsetT) - 1))]; 96 #endif 97 98 public: 99 100 /// Constructor DiscardOutputIterator(OffsetT offset=0)101 __host__ __device__ __forceinline__ DiscardOutputIterator( 102 OffsetT offset = 0) ///< Base offset 103 : 104 offset(offset) 105 {} 106 107 /// Postfix increment operator ++(int)108 __host__ __device__ __forceinline__ self_type operator++(int) 109 { 110 self_type retval = *this; 111 offset++; 112 return retval; 113 } 114 115 /// Prefix increment operator ++()116 __host__ __device__ __forceinline__ self_type operator++() 117 { 118 offset++; 119 return *this; 120 } 121 122 /// Indirection operator *()123 __host__ __device__ __forceinline__ self_type& operator*() 124 { 125 // return self reference, which can be assigned to anything 126 return *this; 127 } 128 129 /// Addition 130 template <typename Distance> operator +(Distance n) const131 __host__ __device__ __forceinline__ self_type operator+(Distance n) const 132 { 133 self_type retval(offset + n); 134 return retval; 135 } 136 137 /// Addition assignment 138 template <typename Distance> operator +=(Distance n)139 __host__ __device__ __forceinline__ self_type& operator+=(Distance n) 140 { 141 offset += n; 142 return *this; 143 } 144 145 /// Subtraction 146 template <typename Distance> operator -(Distance n) const147 __host__ __device__ __forceinline__ self_type operator-(Distance n) const 148 { 149 self_type retval(offset - n); 150 return retval; 151 } 152 153 /// Subtraction assignment 154 template <typename Distance> operator -=(Distance n)155 __host__ __device__ __forceinline__ self_type& operator-=(Distance n) 156 { 157 offset -= n; 158 return *this; 159 } 160 161 /// Distance operator -(self_type other) const162 __host__ __device__ __forceinline__ difference_type operator-(self_type other) const 163 { 164 return offset - other.offset; 165 } 166 167 /// Array subscript 168 template <typename Distance> operator [](Distance n)169 __host__ __device__ __forceinline__ self_type& operator[](Distance n) 170 { 171 // return self reference, which can be assigned to anything 172 return *this; 173 } 174 175 /// Structure dereference operator ->()176 __host__ __device__ __forceinline__ pointer operator->() 177 { 178 return; 179 } 180 181 /// Assignment to self (no-op) operator =(self_type const & other)182 __host__ __device__ __forceinline__ void operator=(self_type const& other) 183 { 184 offset = other.offset; 185 } 186 187 /// Assignment to anything else (no-op) 188 template<typename T> operator =(T const &)189 __host__ __device__ __forceinline__ void operator=(T const&) 190 {} 191 192 /// Cast to void* operator operator void*() const193 __host__ __device__ __forceinline__ operator void*() const { return NULL; } 194 195 /// Equal to operator ==(const self_type & rhs)196 __host__ __device__ __forceinline__ bool operator==(const self_type& rhs) 197 { 198 return (offset == rhs.offset); 199 } 200 201 /// Not equal to operator !=(const self_type & rhs)202 __host__ __device__ __forceinline__ bool operator!=(const self_type& rhs) 203 { 204 return (offset != rhs.offset); 205 } 206 207 /// ostream operator operator <<(std::ostream & os,const self_type & itr)208 friend std::ostream& operator<<(std::ostream& os, const self_type& itr) 209 { 210 os << "[" << itr.offset << "]"; 211 return os; 212 } 213 214 }; 215 216 217 /** @} */ // end group UtilIterator 218 219 } // CUB namespace 220 CUB_NS_POSTFIX // Optional outer namespace(s) 221