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