/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/iterator/cache_modified_output_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_output_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_output_iterator.hpp Source File
cache_modified_output_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_OUTPUT_ITERATOR_HPP_
31 #define HIPCUB_ROCPRIM_ITERATOR_CACHE_MODIFIED_OUTPUT_ITERATOR_HPP_
32 
33 #include <iterator>
34 #include <iostream>
35 
36 #include "../thread/thread_load.hpp"
37 #include "../thread/thread_store.hpp"
38 #include "../util_type.hpp"
39 
40 #if (THRUST_VERSION >= 100700)
41  // This iterator is compatible with Thrust API 1.7 and newer
42  #include <thrust/iterator/iterator_facade.h>
43  #include <thrust/iterator/iterator_traits.h>
44 #endif // THRUST_VERSION
45 
46 
47 BEGIN_HIPCUB_NAMESPACE
48 
49 template <
50  CacheStoreModifier MODIFIER,
51  typename ValueType,
52  typename OffsetT = ptrdiff_t>
54 {
55 private:
56 
57  // Proxy object
58  struct Reference
59  {
60  ValueType* ptr;
61 
63  __host__ __device__ __forceinline__ Reference(ValueType* ptr) : ptr(ptr) {}
64 
66  __device__ __forceinline__ ValueType operator =(ValueType val)
67  {
68  ThreadStore<MODIFIER>(ptr, val);
69  return val;
70  }
71  };
72 
73 public:
74 
75  // Required iterator traits
77  typedef OffsetT difference_type;
78  typedef void value_type;
79  typedef void pointer;
80  typedef Reference reference;
81  typedef std::random_access_iterator_tag iterator_category;
82 
83 private:
84 
85  ValueType* ptr;
86 
87 public:
88 
90  template <typename QualifiedValueType>
91  __host__ __device__ __forceinline__ CacheModifiedOutputIterator(
92  QualifiedValueType* ptr)
93  :
94  ptr(const_cast<typename std::remove_cv<QualifiedValueType>::type *>(ptr))
95  {}
96 
98  __host__ __device__ __forceinline__ self_type operator++(int)
99  {
100  self_type retval = *this;
101  ptr++;
102  return retval;
103  }
104 
105 
107  __host__ __device__ __forceinline__ self_type operator++()
108  {
109  ptr++;
110  return *this;
111  }
112 
114  __host__ __device__ __forceinline__ reference operator*() const
115  {
116  return Reference(ptr);
117  }
118 
120  template <typename Distance>
121  __host__ __device__ __forceinline__ self_type operator+(Distance n) const
122  {
123  self_type retval(ptr + n);
124  return retval;
125  }
126 
128  template <typename Distance>
129  __host__ __device__ __forceinline__ self_type& operator+=(Distance n)
130  {
131  ptr += n;
132  return *this;
133  }
134 
136  template <typename Distance>
137  __host__ __device__ __forceinline__ self_type operator-(Distance n) const
138  {
139  self_type retval(ptr - n);
140  return retval;
141  }
142 
144  template <typename Distance>
145  __host__ __device__ __forceinline__ self_type& operator-=(Distance n)
146  {
147  ptr -= n;
148  return *this;
149  }
150 
152  __host__ __device__ __forceinline__ difference_type operator-(self_type other) const
153  {
154  return ptr - other.ptr;
155  }
156 
158  template <typename Distance>
159  __host__ __device__ __forceinline__ reference operator[](Distance n) const
160  {
161  return Reference(ptr + n);
162  }
163 
165  __host__ __device__ __forceinline__ bool operator==(const self_type& rhs)
166  {
167  return (ptr == rhs.ptr);
168  }
169 
171  __host__ __device__ __forceinline__ bool operator!=(const self_type& rhs)
172  {
173  return (ptr != rhs.ptr);
174  }
175 
176 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
177 
179  friend std::ostream& operator<<(std::ostream& os, const self_type& itr)
180  {
181  (void)itr;
182  return os;
183  }
184 
185 #endif
186 };
187 
188 END_HIPCUB_NAMESPACE
189 
190 #endif
Definition: cache_modified_output_iterator.hpp:54
CacheModifiedOutputIterator self_type
My own type.
Definition: cache_modified_output_iterator.hpp:76
__host__ __device__ __forceinline__ bool operator!=(const self_type &rhs)
Not equal to.
Definition: cache_modified_output_iterator.hpp:171
__host__ __device__ __forceinline__ self_type operator++(int)
Postfix increment.
Definition: cache_modified_output_iterator.hpp:98
__host__ __device__ __forceinline__ reference operator[](Distance n) const
Array subscript.
Definition: cache_modified_output_iterator.hpp:159
__host__ __device__ __forceinline__ self_type & operator+=(Distance n)
Addition assignment.
Definition: cache_modified_output_iterator.hpp:129
Reference reference
The type of a reference to an element the iterator can point to.
Definition: cache_modified_output_iterator.hpp:80
__host__ __device__ __forceinline__ reference operator*() const
Indirection.
Definition: cache_modified_output_iterator.hpp:114
void pointer
The type of a pointer to an element the iterator can point to.
Definition: cache_modified_output_iterator.hpp:79
__host__ __device__ __forceinline__ self_type operator++()
Prefix increment.
Definition: cache_modified_output_iterator.hpp:107
std::random_access_iterator_tag iterator_category
The iterator category.
Definition: cache_modified_output_iterator.hpp:81
__host__ __device__ __forceinline__ CacheModifiedOutputIterator(QualifiedValueType *ptr)
Constructor.
Definition: cache_modified_output_iterator.hpp:91
OffsetT difference_type
Type to express the result of subtracting one iterator from another.
Definition: cache_modified_output_iterator.hpp:77
__host__ __device__ __forceinline__ bool operator==(const self_type &rhs)
Equal to.
Definition: cache_modified_output_iterator.hpp:165
__host__ __device__ __forceinline__ self_type operator+(Distance n) const
Addition.
Definition: cache_modified_output_iterator.hpp:121
__host__ __device__ __forceinline__ self_type operator-(Distance n) const
Subtraction.
Definition: cache_modified_output_iterator.hpp:137
__host__ __device__ __forceinline__ difference_type operator-(self_type other) const
Distance.
Definition: cache_modified_output_iterator.hpp:152
__host__ __device__ __forceinline__ self_type & operator-=(Distance n)
Subtraction assignment.
Definition: cache_modified_output_iterator.hpp:145
void value_type
The type of the element the iterator can point to.
Definition: cache_modified_output_iterator.hpp:78