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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/device/device_segmented_sort.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/device/device_segmented_sort.hpp Source File
device_segmented_sort.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_DEVICE_DEVICE_SEGMENTED_SORT_HPP_
31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_SEGMENTED_SORT_HPP_
32 
33 #include "../../../config.hpp"
34 
35 #include "../util_type.hpp"
36 
37 #include <rocprim/device/device_segmented_radix_sort.hpp>
38 
39 BEGIN_HIPCUB_NAMESPACE
40 
41 struct DeviceSegmentedSort
42 {
43  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
44  HIPCUB_RUNTIME_FUNCTION static
45  hipError_t SortPairs(void * d_temp_storage,
46  size_t& temp_storage_bytes,
47  const KeyT * d_keys_in,
48  KeyT * d_keys_out,
49  const ValueT * d_values_in,
50  ValueT * d_values_out,
51  int num_items,
52  int num_segments,
53  OffsetIteratorT d_begin_offsets,
54  OffsetIteratorT d_end_offsets,
55  hipStream_t stream = 0,
56  bool debug_synchronous = false)
57  {
58  return ::rocprim::segmented_radix_sort_pairs(
59  d_temp_storage, temp_storage_bytes,
60  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,
61  num_segments, d_begin_offsets, d_end_offsets,
62  0, sizeof(KeyT) * 8,
63  stream, debug_synchronous
64  );
65  }
66 
67  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
68  HIPCUB_RUNTIME_FUNCTION static
69  hipError_t SortPairs(void * d_temp_storage,
70  size_t& temp_storage_bytes,
71  DoubleBuffer<KeyT>& d_keys,
72  DoubleBuffer<ValueT>& d_values,
73  int num_items,
74  int num_segments,
75  OffsetIteratorT d_begin_offsets,
76  OffsetIteratorT d_end_offsets,
77  hipStream_t stream = 0,
78  bool debug_synchronous = false)
79  {
80  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
81  ::rocprim::double_buffer<ValueT> d_values_db = detail::to_double_buffer(d_values);
82  hipError_t error = ::rocprim::segmented_radix_sort_pairs(
83  d_temp_storage, temp_storage_bytes,
84  d_keys_db, d_values_db, num_items,
85  num_segments, d_begin_offsets, d_end_offsets,
86  0, sizeof(KeyT) * 8,
87  stream, debug_synchronous
88  );
89  detail::update_double_buffer(d_keys, d_keys_db);
90  detail::update_double_buffer(d_values, d_values_db);
91  return error;
92  }
93 
94  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
95  HIPCUB_RUNTIME_FUNCTION static
96  hipError_t SortPairsDescending(void * d_temp_storage,
97  size_t& temp_storage_bytes,
98  const KeyT * d_keys_in,
99  KeyT * d_keys_out,
100  const ValueT * d_values_in,
101  ValueT * d_values_out,
102  int num_items,
103  int num_segments,
104  OffsetIteratorT d_begin_offsets,
105  OffsetIteratorT d_end_offsets,
106  hipStream_t stream = 0,
107  bool debug_synchronous = false)
108  {
109  return ::rocprim::segmented_radix_sort_pairs_desc(
110  d_temp_storage, temp_storage_bytes,
111  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,
112  num_segments, d_begin_offsets, d_end_offsets,
113  0, sizeof(KeyT) * 8,
114  stream, debug_synchronous
115  );
116  }
117 
118  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
119  HIPCUB_RUNTIME_FUNCTION static
120  hipError_t SortPairsDescending(void * d_temp_storage,
121  size_t& temp_storage_bytes,
122  DoubleBuffer<KeyT>& d_keys,
123  DoubleBuffer<ValueT>& d_values,
124  int num_items,
125  int num_segments,
126  OffsetIteratorT d_begin_offsets,
127  OffsetIteratorT d_end_offsets,
128  hipStream_t stream = 0,
129  bool debug_synchronous = false)
130  {
131  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
132  ::rocprim::double_buffer<ValueT> d_values_db = detail::to_double_buffer(d_values);
133  hipError_t error = ::rocprim::segmented_radix_sort_pairs_desc(
134  d_temp_storage, temp_storage_bytes,
135  d_keys_db, d_values_db, num_items,
136  num_segments, d_begin_offsets, d_end_offsets,
137  0, sizeof(KeyT) * 8,
138  stream, debug_synchronous
139  );
140  detail::update_double_buffer(d_keys, d_keys_db);
141  detail::update_double_buffer(d_values, d_values_db);
142  return error;
143  }
144 
145  template<typename KeyT, typename OffsetIteratorT>
146  HIPCUB_RUNTIME_FUNCTION static
147  hipError_t SortKeys(void * d_temp_storage,
148  size_t& temp_storage_bytes,
149  const KeyT * d_keys_in,
150  KeyT * d_keys_out,
151  int num_items,
152  int num_segments,
153  OffsetIteratorT d_begin_offsets,
154  OffsetIteratorT d_end_offsets,
155  hipStream_t stream = 0,
156  bool debug_synchronous = false)
157  {
158  return ::rocprim::segmented_radix_sort_keys(
159  d_temp_storage, temp_storage_bytes,
160  d_keys_in, d_keys_out, num_items,
161  num_segments, d_begin_offsets, d_end_offsets,
162  0, sizeof(KeyT) * 8,
163  stream, debug_synchronous
164  );
165  }
166 
167  template<typename KeyT, typename OffsetIteratorT>
168  HIPCUB_RUNTIME_FUNCTION static
169  hipError_t SortKeys(void * d_temp_storage,
170  size_t& temp_storage_bytes,
171  DoubleBuffer<KeyT>& d_keys,
172  int num_items,
173  int num_segments,
174  OffsetIteratorT d_begin_offsets,
175  OffsetIteratorT d_end_offsets,
176  hipStream_t stream = 0,
177  bool debug_synchronous = false)
178  {
179  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
180  hipError_t error = ::rocprim::segmented_radix_sort_keys(
181  d_temp_storage, temp_storage_bytes,
182  d_keys_db, num_items,
183  num_segments, d_begin_offsets, d_end_offsets,
184  0, sizeof(KeyT) * 8,
185  stream, debug_synchronous
186  );
187  detail::update_double_buffer(d_keys, d_keys_db);
188  return error;
189  }
190 
191  template<typename KeyT, typename OffsetIteratorT>
192  HIPCUB_RUNTIME_FUNCTION static
193  hipError_t SortKeysDescending(void * d_temp_storage,
194  size_t& temp_storage_bytes,
195  const KeyT * d_keys_in,
196  KeyT * d_keys_out,
197  int num_items,
198  int num_segments,
199  OffsetIteratorT d_begin_offsets,
200  OffsetIteratorT d_end_offsets,
201  hipStream_t stream = 0,
202  bool debug_synchronous = false)
203  {
204  return ::rocprim::segmented_radix_sort_keys_desc(
205  d_temp_storage, temp_storage_bytes,
206  d_keys_in, d_keys_out, num_items,
207  num_segments, d_begin_offsets, d_end_offsets,
208  0, sizeof(KeyT) * 8,
209  stream, debug_synchronous
210  );
211  }
212 
213  template<typename KeyT, typename OffsetIteratorT>
214  HIPCUB_RUNTIME_FUNCTION static
215  hipError_t SortKeysDescending(void * d_temp_storage,
216  size_t& temp_storage_bytes,
217  DoubleBuffer<KeyT>& d_keys,
218  int num_items,
219  int num_segments,
220  OffsetIteratorT d_begin_offsets,
221  OffsetIteratorT d_end_offsets,
222  hipStream_t stream = 0,
223  bool debug_synchronous = false)
224  {
225  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
226  hipError_t error = ::rocprim::segmented_radix_sort_keys_desc(
227  d_temp_storage, temp_storage_bytes,
228  d_keys_db, num_items,
229  num_segments, d_begin_offsets, d_end_offsets,
230  0, sizeof(KeyT) * 8,
231  stream, debug_synchronous
232  );
233  detail::update_double_buffer(d_keys, d_keys_db);
234  return error;
235  }
236 
237  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
238  HIPCUB_RUNTIME_FUNCTION static
239  hipError_t StableSortPairs(void * d_temp_storage,
240  size_t& temp_storage_bytes,
241  const KeyT * d_keys_in,
242  KeyT * d_keys_out,
243  const ValueT * d_values_in,
244  ValueT * d_values_out,
245  int num_items,
246  int num_segments,
247  OffsetIteratorT d_begin_offsets,
248  OffsetIteratorT d_end_offsets,
249  hipStream_t stream = 0,
250  bool debug_synchronous = false)
251  {
252  return SortPairs(
253  d_temp_storage, temp_storage_bytes,
254  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,
255  num_segments, d_begin_offsets, d_end_offsets,
256  stream, debug_synchronous
257  );
258  }
259 
260  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
261  HIPCUB_RUNTIME_FUNCTION static
262  hipError_t StableSortPairs(void * d_temp_storage,
263  size_t& temp_storage_bytes,
264  DoubleBuffer<KeyT>& d_keys,
265  DoubleBuffer<ValueT>& d_values,
266  int num_items,
267  int num_segments,
268  OffsetIteratorT d_begin_offsets,
269  OffsetIteratorT d_end_offsets,
270  hipStream_t stream = 0,
271  bool debug_synchronous = false)
272  {
273  return SortPairs(
274  d_temp_storage, temp_storage_bytes,
275  d_keys, d_values, num_items,
276  num_segments, d_begin_offsets, d_end_offsets,
277  stream, debug_synchronous
278  );
279  }
280 
281  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
282  HIPCUB_RUNTIME_FUNCTION static
283  hipError_t StableSortPairsDescending(void * d_temp_storage,
284  size_t& temp_storage_bytes,
285  const KeyT * d_keys_in,
286  KeyT * d_keys_out,
287  const ValueT * d_values_in,
288  ValueT * d_values_out,
289  int num_items,
290  int num_segments,
291  OffsetIteratorT d_begin_offsets,
292  OffsetIteratorT d_end_offsets,
293  hipStream_t stream = 0,
294  bool debug_synchronous = false)
295  {
296  return SortPairsDescending(
297  d_temp_storage, temp_storage_bytes,
298  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,
299  num_segments, d_begin_offsets, d_end_offsets,
300  stream, debug_synchronous
301  );
302  }
303 
304  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
305  HIPCUB_RUNTIME_FUNCTION static
306  hipError_t StableSortPairsDescending(void * d_temp_storage,
307  size_t& temp_storage_bytes,
308  DoubleBuffer<KeyT>& d_keys,
309  DoubleBuffer<ValueT>& d_values,
310  int num_items,
311  int num_segments,
312  OffsetIteratorT d_begin_offsets,
313  OffsetIteratorT d_end_offsets,
314  hipStream_t stream = 0,
315  bool debug_synchronous = false)
316  {
317  return SortPairsDescending(
318  d_temp_storage, temp_storage_bytes,
319  d_keys, d_values, num_items,
320  num_segments, d_begin_offsets, d_end_offsets,
321  stream, debug_synchronous
322  );
323  }
324 
325  template<typename KeyT, typename OffsetIteratorT>
326  HIPCUB_RUNTIME_FUNCTION static
327  hipError_t StableSortKeys(void * d_temp_storage,
328  size_t& temp_storage_bytes,
329  const KeyT * d_keys_in,
330  KeyT * d_keys_out,
331  int num_items,
332  int num_segments,
333  OffsetIteratorT d_begin_offsets,
334  OffsetIteratorT d_end_offsets,
335  hipStream_t stream = 0,
336  bool debug_synchronous = false)
337  {
338  return SortKeys(
339  d_temp_storage, temp_storage_bytes,
340  d_keys_in, d_keys_out, num_items,
341  num_segments, d_begin_offsets, d_end_offsets,
342  stream, debug_synchronous
343  );
344  }
345 
346  template<typename KeyT, typename OffsetIteratorT>
347  HIPCUB_RUNTIME_FUNCTION static
348  hipError_t StableSortKeys(void * d_temp_storage,
349  size_t& temp_storage_bytes,
350  DoubleBuffer<KeyT>& d_keys,
351  int num_items,
352  int num_segments,
353  OffsetIteratorT d_begin_offsets,
354  OffsetIteratorT d_end_offsets,
355  hipStream_t stream = 0,
356  bool debug_synchronous = false)
357  {
358  return SortKeys(
359  d_temp_storage, temp_storage_bytes,
360  d_keys, num_items,
361  num_segments, d_begin_offsets, d_end_offsets,
362  stream, debug_synchronous
363  );
364  }
365 
366  template<typename KeyT, typename OffsetIteratorT>
367  HIPCUB_RUNTIME_FUNCTION static
368  hipError_t StableSortKeysDescending(void * d_temp_storage,
369  size_t& temp_storage_bytes,
370  const KeyT * d_keys_in,
371  KeyT * d_keys_out,
372  int num_items,
373  int num_segments,
374  OffsetIteratorT d_begin_offsets,
375  OffsetIteratorT d_end_offsets,
376  hipStream_t stream = 0,
377  bool debug_synchronous = false)
378  {
379  return SortKeysDescending(
380  d_temp_storage, temp_storage_bytes,
381  d_keys_in, d_keys_out, num_items,
382  num_segments, d_begin_offsets, d_end_offsets,
383  stream, debug_synchronous
384  );
385  }
386 
387  template<typename KeyT, typename OffsetIteratorT>
388  HIPCUB_RUNTIME_FUNCTION static
389  hipError_t StableSortKeysDescending(void * d_temp_storage,
390  size_t& temp_storage_bytes,
391  DoubleBuffer<KeyT>& d_keys,
392  int num_items,
393  int num_segments,
394  OffsetIteratorT d_begin_offsets,
395  OffsetIteratorT d_end_offsets,
396  hipStream_t stream = 0,
397  bool debug_synchronous = false)
398  {
399  return SortKeysDescending(
400  d_temp_storage, temp_storage_bytes,
401  d_keys, num_items,
402  num_segments, d_begin_offsets, d_end_offsets,
403  stream, debug_synchronous
404  );
405  }
406 };
407 
408 END_HIPCUB_NAMESPACE
409 
410 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_SEGMENTED_SORT_HPP_