CUB
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups
cache_modified_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 #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 
52 CUB_NS_PREFIX
53 
55 namespace cub {
56 
57 
58 
103 template <
104  CacheLoadModifier MODIFIER,
105  typename ValueType,
106  typename OffsetT = ptrdiff_t>
108 {
109 public:
110 
111  // Required iterator traits
113  typedef OffsetT difference_type;
114  typedef ValueType value_type;
115  typedef ValueType* pointer;
116  typedef ValueType reference;
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;
126 #else
127  typedef std::random_access_iterator_tag iterator_category;
128 #endif // THRUST_VERSION
129 
130 
131 public:
132 
134  ValueType* ptr;
135 
137  template <typename QualifiedValueType>
138  __host__ __device__ __forceinline__ CacheModifiedInputIterator(
139  QualifiedValueType* ptr)
140  :
141  ptr(const_cast<typename RemoveQualifiers<QualifiedValueType>::Type *>(ptr))
142  {}
143 
145  __host__ __device__ __forceinline__ self_type operator++(int)
146  {
147  self_type retval = *this;
148  ptr++;
149  return retval;
150  }
151 
153  __host__ __device__ __forceinline__ self_type operator++()
154  {
155  ptr++;
156  return *this;
157  }
158 
160  __device__ __forceinline__ reference operator*() const
161  {
162  return ThreadLoad<MODIFIER>(ptr);
163  }
164 
166  template <typename Distance>
167  __host__ __device__ __forceinline__ self_type operator+(Distance n) const
168  {
169  self_type retval(ptr + n);
170  return retval;
171  }
172 
174  template <typename Distance>
175  __host__ __device__ __forceinline__ self_type& operator+=(Distance n)
176  {
177  ptr += n;
178  return *this;
179  }
180 
182  template <typename Distance>
183  __host__ __device__ __forceinline__ self_type operator-(Distance n) const
184  {
185  self_type retval(ptr - n);
186  return retval;
187  }
188 
190  template <typename Distance>
191  __host__ __device__ __forceinline__ self_type& operator-=(Distance n)
192  {
193  ptr -= n;
194  return *this;
195  }
196 
198  __host__ __device__ __forceinline__ difference_type operator-(self_type other) const
199  {
200  return ptr - other.ptr;
201  }
202 
204  template <typename Distance>
205  __device__ __forceinline__ reference operator[](Distance n) const
206  {
207  return ThreadLoad<MODIFIER>(ptr + n);
208  }
209 
211  __device__ __forceinline__ pointer operator->()
212  {
213  return &ThreadLoad<MODIFIER>(ptr);
214  }
215 
217  __host__ __device__ __forceinline__ bool operator==(const self_type& rhs)
218  {
219  return (ptr == rhs.ptr);
220  }
221 
223  __host__ __device__ __forceinline__ bool operator!=(const self_type& rhs)
224  {
225  return (ptr != rhs.ptr);
226  }
227 
229  friend std::ostream& operator<<(std::ostream& os, const self_type& itr)
230  {
231  return os;
232  }
233 };
234 
235 
236  // end group UtilIterator
238 
239 } // CUB namespace
240 CUB_NS_POSTFIX // Optional outer namespace(s)