/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.4.4/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.4.4/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, 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 = CubDebug(hipMemcpyAsync(d_counters, counters, sizeof(OffsetT) * 2, hipMemcpyHostToDevice, stream));
136  return result;
137  }
138 
140  HIPCUB_DEVICE hipError_t ResetDrain(hipStream_t stream = 0)
141  {
142  hipError_t result = hipErrorUnknown;
143  (void)stream;
144  d_counters[DRAIN] = 0;
145  result = hipSuccess;
146  return result;
147  }
148 
149  HIPCUB_HOST hipError_t ResetDrain(hipStream_t stream = 0)
150  {
151  hipError_t result = hipErrorUnknown;
152  result = CubDebug(hipMemsetAsync(d_counters + DRAIN, 0, sizeof(OffsetT), stream));
153  return result;
154  }
155 
156 
158  HIPCUB_DEVICE hipError_t ResetFill(hipStream_t stream = 0)
159  {
160  hipError_t result = hipErrorUnknown;
161  (void)stream;
162  d_counters[FILL] = 0;
163  result = hipSuccess;
164  return result;
165  }
166 
167  HIPCUB_HOST hipError_t ResetFill(hipStream_t stream = 0)
168  {
169  hipError_t result = hipErrorUnknown;
170  result = CubDebug(hipMemsetAsync(d_counters + FILL, 0, sizeof(OffsetT), stream));
171  return result;
172  }
173 
174 
176  HIPCUB_DEVICE hipError_t FillSize(
177  OffsetT &fill_size,
178  hipStream_t stream = 0)
179  {
180  hipError_t result = hipErrorUnknown;
181  (void)stream;
182  fill_size = d_counters[FILL];
183  result = hipSuccess;
184  return result;
185  }
186 
187  HIPCUB_HOST hipError_t FillSize(
188  OffsetT &fill_size,
189  hipStream_t stream = 0)
190  {
191  hipError_t result = hipErrorUnknown;
192  result = CubDebug(hipMemcpyAsync(&fill_size, d_counters + FILL, sizeof(OffsetT), hipMemcpyDeviceToHost, stream));
193  return result;
194  }
195 
196 
198  HIPCUB_DEVICE OffsetT Drain(OffsetT num_items)
199  {
200  return atomicAdd(d_counters + DRAIN, num_items);
201  }
202 
203 
205  HIPCUB_DEVICE OffsetT Fill(OffsetT num_items)
206  {
207  return atomicAdd(d_counters + FILL, num_items);
208  }
209 };
210 
211 
212 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
213 
214 
218 template <typename OffsetT>
219 __global__ void FillAndResetDrainKernel(
220  GridQueue<OffsetT> grid_queue,
221  OffsetT num_items)
222 {
223  grid_queue.FillAndResetDrain(num_items);
224 }
225 
226 
227 
228 #endif // DOXYGEN_SHOULD_SKIP_THIS
229 
230  // end group GridModule
232 
233 END_HIPCUB_NAMESPACE
234 
235 #endif // HIPCUB_ROCPRIM_GRID_GRID_QUEUE_HPP_
hipcub::GridQueue
GridQueue is a descriptor utility for dynamic queue management.
Definition: grid_queue.hpp:102