/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/iterator/cache_modified_input_iterator.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/iterator/cache_modified_input_iterator.hpp Source File#

hipCUB: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/iterator/cache_modified_input_iterator.hpp Source File
cache_modified_input_iterator.hpp
1 /******************************************************************************
2  * Copyright (c) 2010-2011, Duane Merrill. All rights reserved.
3  * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
4  * Modifications Copyright (c) 2017-2020, Advanced Micro Devices, Inc. All rights reserved.
5  *
6  * Redistribution and use in source and binary forms, with or without
7  * modification, are permitted provided that the following conditions are met:
8  * * Redistributions of source code must retain the above copyright
9  * notice, this list of conditions and the following disclaimer.
10  * * Redistributions in binary form must reproduce the above copyright
11  * notice, this list of conditions and the following disclaimer in the
12  * documentation and/or other materials provided with the distribution.
13  * * Neither the name of the NVIDIA CORPORATION nor the
14  * names of its contributors may be used to endorse or promote products
15  * derived from this software without specific prior written permission.
16  *
17  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
18  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
19  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
21  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
22  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
23  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
24  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
26  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27  *
28  ******************************************************************************/
29 
30 #ifndef HIPCUB_ROCPRIM_ITERATOR_CACHE_MODIFIED_INPUT_ITERATOR_HPP_
31 #define HIPCUB_ROCPRIM_ITERATOR_CACHE_MODIFIED_INPUT_ITERATOR_HPP_
32 
33 #include <iterator>
34 #include <iostream>
35 
36 #include "../thread/thread_load.hpp"
37 #include "../util_type.hpp"
38 
39 #if (THRUST_VERSION >= 100700)
40  // This iterator is compatible with Thrust API 1.7 and newer
41  #include <thrust/iterator/iterator_facade.h>
42  #include <thrust/iterator/iterator_traits.h>
43 #endif // THRUST_VERSION
44 
45 BEGIN_HIPCUB_NAMESPACE
46 
47 template <
48  CacheLoadModifier MODIFIER,
49  typename ValueType,
50  typename OffsetT = ptrdiff_t>
52 {
53 public:
54 
55  // Required iterator traits
57  typedef OffsetT difference_type;
58  typedef ValueType value_type;
59  typedef ValueType* pointer;
60  typedef ValueType reference;
61  typedef std::random_access_iterator_tag iterator_category;
62 
63 public:
64 
66  ValueType* ptr;
67 
69  __host__ __device__ __forceinline__ CacheModifiedInputIterator(
70  ValueType* ptr)
71  :
72  ptr(const_cast<typename std::remove_cv<ValueType>::type *>(ptr))
73  {}
74 
76  __host__ __device__ __forceinline__ self_type operator++(int)
77  {
78  self_type retval = *this;
79  ptr++;
80  return retval;
81  }
82 
84  __host__ __device__ __forceinline__ self_type operator++()
85  {
86  ptr++;
87  return *this;
88  }
89 
91  __device__ __forceinline__ reference operator*() const
92  {
93  return ThreadLoad<MODIFIER>(ptr);
94  }
95 
97  template <typename Distance>
98  __host__ __device__ __forceinline__ self_type operator+(Distance n) const
99  {
100  self_type retval(ptr + n);
101  return retval;
102  }
103 
105  template <typename Distance>
106  __host__ __device__ __forceinline__ self_type& operator+=(Distance n)
107  {
108  ptr += n;
109  return *this;
110  }
111 
113  template <typename Distance>
114  __host__ __device__ __forceinline__ self_type operator-(Distance n) const
115  {
116  self_type retval(ptr - n);
117  return retval;
118  }
119 
121  template <typename Distance>
122  __host__ __device__ __forceinline__ self_type& operator-=(Distance n)
123  {
124  ptr -= n;
125  return *this;
126  }
127 
129  __host__ __device__ __forceinline__ difference_type operator-(self_type other) const
130  {
131  return ptr - other.ptr;
132  }
133 
135  template <typename Distance>
136  __device__ __forceinline__ reference operator[](Distance n) const
137  {
138  return ThreadLoad<MODIFIER>(ptr + n);
139  }
140 
142  __device__ __forceinline__ pointer operator->()
143  {
144  return &ThreadLoad<MODIFIER>(ptr);
145  }
146 
148  __host__ __device__ __forceinline__ bool operator==(const self_type& rhs)
149  {
150  return (ptr == rhs.ptr);
151  }
152 
154  __host__ __device__ __forceinline__ bool operator!=(const self_type& rhs)
155  {
156  return (ptr != rhs.ptr);
157  }
158 
159 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
160 
162  friend std::ostream& operator<<(std::ostream& os, const self_type& /*itr*/)
163  {
164  return os;
165  }
166 
167 #endif
168 
169 };
170 
171 END_HIPCUB_NAMESPACE
172 
173 #endif // HIPCUB_ROCPRIM_ITERATOR_CACHE_MODIFIED_INPUT_ITERATOR_HPP_
Definition: cache_modified_input_iterator.hpp:52
__host__ __device__ __forceinline__ self_type operator++()
Prefix increment.
Definition: cache_modified_input_iterator.hpp:84
__host__ __device__ __forceinline__ self_type operator-(Distance n) const
Subtraction.
Definition: cache_modified_input_iterator.hpp:114
OffsetT difference_type
Type to express the result of subtracting one iterator from another.
Definition: cache_modified_input_iterator.hpp:57
ValueType reference
The type of a reference to an element the iterator can point to.
Definition: cache_modified_input_iterator.hpp:60
__device__ __forceinline__ reference operator[](Distance n) const
Array subscript.
Definition: cache_modified_input_iterator.hpp:136
CacheModifiedInputIterator self_type
My own type.
Definition: cache_modified_input_iterator.hpp:56
ValueType * pointer
The type of a pointer to an element the iterator can point to.
Definition: cache_modified_input_iterator.hpp:59
__host__ __device__ __forceinline__ bool operator==(const self_type &rhs)
Equal to.
Definition: cache_modified_input_iterator.hpp:148
__device__ __forceinline__ reference operator*() const
Indirection.
Definition: cache_modified_input_iterator.hpp:91
ValueType * ptr
Wrapped native pointer.
Definition: cache_modified_input_iterator.hpp:66
__host__ __device__ __forceinline__ self_type operator+(Distance n) const
Addition.
Definition: cache_modified_input_iterator.hpp:98
__host__ __device__ __forceinline__ difference_type operator-(self_type other) const
Distance.
Definition: cache_modified_input_iterator.hpp:129
ValueType value_type
The type of the element the iterator can point to.
Definition: cache_modified_input_iterator.hpp:58
__host__ __device__ __forceinline__ bool operator!=(const self_type &rhs)
Not equal to.
Definition: cache_modified_input_iterator.hpp:154
std::random_access_iterator_tag iterator_category
The iterator category.
Definition: cache_modified_input_iterator.hpp:61
__device__ __forceinline__ pointer operator->()
Structure dereference.
Definition: cache_modified_input_iterator.hpp:142
__host__ __device__ __forceinline__ self_type operator++(int)
Postfix increment.
Definition: cache_modified_input_iterator.hpp:76
__host__ __device__ __forceinline__ CacheModifiedInputIterator(ValueType *ptr)
Constructor.
Definition: cache_modified_input_iterator.hpp:69
__host__ __device__ __forceinline__ self_type & operator+=(Distance n)
Addition assignment.
Definition: cache_modified_input_iterator.hpp:106
__host__ __device__ __forceinline__ self_type & operator-=(Distance n)
Subtraction assignment.
Definition: cache_modified_input_iterator.hpp:122