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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/grid/grid_queue.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/grid/grid_queue.hpp Source File
grid_queue.hpp
1 /******************************************************************************
2  * Copyright (c) 2011, Duane Merrill. All rights reserved.
3  * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
4  * Modifications Copyright (c) 2021-2023, 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_GRID_GRID_QUEUE_HPP_
31 #define HIPCUB_ROCPRIM_GRID_GRID_QUEUE_HPP_
32 
33 #include <type_traits>
34 
35 #include "../../../config.hpp"
36 
37 BEGIN_HIPCUB_NAMESPACE
38 
74 template <typename OffsetT>
75 class GridQueue
76 {
77 private:
78 
80  enum
81  {
82  FILL = 0,
83  DRAIN = 1,
84  };
85 
87  OffsetT *d_counters;
88 
89 public:
90 
92  __host__ __device__ __forceinline__
93  static size_t AllocationSize()
94  {
95  return sizeof(OffsetT) * 2;
96  }
97 
98 
100  __host__ __device__ __forceinline__ GridQueue()
101  :
102  d_counters(NULL)
103  {}
104 
105 
107  __host__ __device__ __forceinline__ GridQueue(
108  void *d_storage)
109  :
110  d_counters((OffsetT*) d_storage)
111  {}
112 
113 
115  HIPCUB_DEVICE hipError_t FillAndResetDrain(
116  OffsetT fill_size,
117  hipStream_t stream = 0)
118  {
119  hipError_t result = hipErrorUnknown;
120  (void)stream;
121  d_counters[FILL] = fill_size;
122  d_counters[DRAIN] = 0;
123  result = hipSuccess;
124  return result;
125  }
126 
127  HIPCUB_HOST hipError_t FillAndResetDrain(
128  OffsetT fill_size,
129  hipStream_t stream = 0)
130  {
131  hipError_t result = hipErrorUnknown;
132  OffsetT counters[2];
133  counters[FILL] = fill_size;
134  counters[DRAIN] = 0;
135  result = HipcubDebug(hipMemcpyAsync(d_counters,
136  counters,
137  sizeof(OffsetT) * 2,
138  hipMemcpyHostToDevice,
139  stream));
140  return result;
141  }
142 
144  HIPCUB_DEVICE hipError_t ResetDrain(hipStream_t stream = 0)
145  {
146  hipError_t result = hipErrorUnknown;
147  (void)stream;
148  d_counters[DRAIN] = 0;
149  result = hipSuccess;
150  return result;
151  }
152 
153  HIPCUB_HOST hipError_t ResetDrain(hipStream_t stream = 0)
154  {
155  hipError_t result = hipErrorUnknown;
156  result = HipcubDebug(hipMemsetAsync(d_counters + DRAIN, 0, sizeof(OffsetT), stream));
157  return result;
158  }
159 
160 
162  HIPCUB_DEVICE hipError_t ResetFill(hipStream_t stream = 0)
163  {
164  hipError_t result = hipErrorUnknown;
165  (void)stream;
166  d_counters[FILL] = 0;
167  result = hipSuccess;
168  return result;
169  }
170 
171  HIPCUB_HOST hipError_t ResetFill(hipStream_t stream = 0)
172  {
173  hipError_t result = hipErrorUnknown;
174  result = HipcubDebug(hipMemsetAsync(d_counters + FILL, 0, sizeof(OffsetT), stream));
175  return result;
176  }
177 
178 
180  HIPCUB_DEVICE hipError_t FillSize(
181  OffsetT &fill_size,
182  hipStream_t stream = 0)
183  {
184  hipError_t result = hipErrorUnknown;
185  (void)stream;
186  fill_size = d_counters[FILL];
187  result = hipSuccess;
188  return result;
189  }
190 
191  HIPCUB_HOST hipError_t FillSize(
192  OffsetT &fill_size,
193  hipStream_t stream = 0)
194  {
195  hipError_t result = hipErrorUnknown;
196  result = HipcubDebug(hipMemcpyAsync(&fill_size,
197  d_counters + FILL,
198  sizeof(OffsetT),
199  hipMemcpyDeviceToHost,
200  stream));
201  return result;
202  }
203 
204 
206  HIPCUB_DEVICE OffsetT Drain(OffsetT num_items)
207  {
208  return atomicAdd(d_counters + DRAIN, num_items);
209  }
210 
211 
213  HIPCUB_DEVICE OffsetT Fill(OffsetT num_items)
214  {
215  return atomicAdd(d_counters + FILL, num_items);
216  }
217 };
218 
219 
220 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
221 
222 
226 template <typename OffsetT>
227 __global__ void FillAndResetDrainKernel(
228  GridQueue<OffsetT> grid_queue,
229  OffsetT num_items)
230 {
231  grid_queue.FillAndResetDrain(num_items);
232 }
233 
234 
235 
236 #endif // DOXYGEN_SHOULD_SKIP_THIS
237 
238  // end group GridModule
240 
241 END_HIPCUB_NAMESPACE
242 
243 #endif // HIPCUB_ROCPRIM_GRID_GRID_QUEUE_HPP_
GridQueue is a descriptor utility for dynamic queue management.
Definition: grid_queue.hpp:76
__device__ hipError_t FillSize(OffsetT &fill_size, hipStream_t stream=0)
Returns the fill-size established by the parent or by the previous kernel.
Definition: grid_queue.hpp:180
__host__ __device__ __forceinline__ GridQueue()
Constructs an invalid GridQueue descriptor.
Definition: grid_queue.hpp:100
__device__ hipError_t ResetDrain(hipStream_t stream=0)
This operation resets the drain so that it may advance to meet the existing fill-size....
Definition: grid_queue.hpp:144
__host__ __device__ static __forceinline__ size_t AllocationSize()
Returns the device allocation size in bytes needed to construct a GridQueue instance.
Definition: grid_queue.hpp:93
__host__ __device__ __forceinline__ GridQueue(void *d_storage)
Constructs a GridQueue descriptor around the device storage allocation.
Definition: grid_queue.hpp:107
__device__ hipError_t FillAndResetDrain(OffsetT fill_size, hipStream_t stream=0)
This operation sets the fill-size and resets the drain counter, preparing the GridQueue for draining ...
Definition: grid_queue.hpp:115
__device__ hipError_t ResetFill(hipStream_t stream=0)
This operation resets the fill counter. To be called by the host or by a kernel prior to that which w...
Definition: grid_queue.hpp:162
__device__ OffsetT Fill(OffsetT num_items)
Fill num_items into the queue. Returns offset from which to write items. To be called from hip kernel...
Definition: grid_queue.hpp:213
__device__ OffsetT Drain(OffsetT num_items)
Drain num_items from the queue. Returns offset from which to read items. To be called from hip kernel...
Definition: grid_queue.hpp:206