CUB
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups
arg_index_input_iterator.cuh
Go to the documentation of this file.
1 /******************************************************************************
2  * Copyright (c) 2011, Duane Merrill. All rights reserved.
3  * Copyright (c) 2011-2016, 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 
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 #include <thrust/version.h>
45 
46 #if (THRUST_VERSION >= 100700)
47  // This iterator is compatible with Thrust API 1.7 and newer
48  #include <thrust/iterator/iterator_facade.h>
49  #include <thrust/iterator/iterator_traits.h>
50 #endif // THRUST_VERSION
51 
53 CUB_NS_PREFIX
54 
56 namespace cub {
57 
108 template <
109  typename InputIteratorT,
110  typename OffsetT = ptrdiff_t>
112 {
113 private:
114 
115  // Data type of input iterator
116  typedef typename std::iterator_traits<InputIteratorT>::value_type T;
117 
118 public:
119 
120 
121  // Required iterator traits
123  typedef OffsetT difference_type;
124  typedef KeyValuePair<difference_type, T> value_type;
125  typedef value_type* pointer;
127 
128 #if (THRUST_VERSION >= 100700)
129  // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods
130  typedef typename thrust::detail::iterator_facade_category<
131  thrust::any_system_tag,
132  thrust::random_access_traversal_tag,
133  value_type,
134  reference
135  >::type iterator_category;
136 #else
137  typedef std::random_access_iterator_tag iterator_category;
138 #endif // THRUST_VERSION
139 
140 private:
141 
142  InputIteratorT itr;
143  difference_type offset;
144 
145 public:
146 
148  __host__ __device__ __forceinline__ ArgIndexInputIterator(
149  InputIteratorT itr,
150  difference_type offset = 0)
151  :
152  itr(itr),
153  offset(offset)
154  {}
155 
157  __host__ __device__ __forceinline__ self_type operator++(int)
158  {
159  self_type retval = *this;
160  offset++;
161  return retval;
162  }
163 
165  __host__ __device__ __forceinline__ self_type operator++()
166  {
167  offset++;
168  return *this;
169  }
170 
172  __host__ __device__ __forceinline__ reference operator*() const
173  {
174  value_type retval;
175  retval.value = itr[offset];
176  retval.key = offset;
177  return retval;
178  }
179 
181  template <typename Distance>
182  __host__ __device__ __forceinline__ self_type operator+(Distance n) const
183  {
184  self_type retval(itr, offset + n);
185  return retval;
186  }
187 
189  template <typename Distance>
190  __host__ __device__ __forceinline__ self_type& operator+=(Distance n)
191  {
192  offset += n;
193  return *this;
194  }
195 
197  template <typename Distance>
198  __host__ __device__ __forceinline__ self_type operator-(Distance n) const
199  {
200  self_type retval(itr, offset - n);
201  return retval;
202  }
203 
205  template <typename Distance>
206  __host__ __device__ __forceinline__ self_type& operator-=(Distance n)
207  {
208  offset -= n;
209  return *this;
210  }
211 
213  __host__ __device__ __forceinline__ difference_type operator-(self_type other) const
214  {
215  return offset - other.offset;
216  }
217 
219  template <typename Distance>
220  __host__ __device__ __forceinline__ reference operator[](Distance n) const
221  {
222  self_type offset = (*this) + n;
223  return *offset;
224  }
225 
227  __host__ __device__ __forceinline__ pointer operator->()
228  {
229  return &(*(*this));
230  }
231 
233  __host__ __device__ __forceinline__ bool operator==(const self_type& rhs)
234  {
235  return ((itr == rhs.itr) && (offset == rhs.offset));
236  }
237 
239  __host__ __device__ __forceinline__ bool operator!=(const self_type& rhs)
240  {
241  return ((itr != rhs.itr) || (offset != rhs.offset));
242  }
243 
245  __host__ __device__ __forceinline__ void normalize()
246  {
247  itr += offset;
248  offset = 0;
249  }
250 
252  friend std::ostream& operator<<(std::ostream& os, const self_type& itr)
253  {
254  return os;
255  }
256 };
257 
258 
259  // end group UtilIterator
261 
262 } // CUB namespace
263 CUB_NS_POSTFIX // Optional outer namespace(s)