/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/device/device_segmented_radix_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_radix_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_radix_sort.hpp Source File
device_segmented_radix_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_RADIX_SORT_HPP_
31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_SEGMENTED_RADIX_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 DeviceSegmentedRadixSort
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  int begin_bit = 0,
56  int end_bit = sizeof(KeyT) * 8,
57  hipStream_t stream = 0,
58  bool debug_synchronous = false)
59  {
60  return ::rocprim::segmented_radix_sort_pairs(
61  d_temp_storage, temp_storage_bytes,
62  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,
63  num_segments, d_begin_offsets, d_end_offsets,
64  begin_bit, end_bit,
65  stream, debug_synchronous
66  );
67  }
68 
69  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
70  HIPCUB_RUNTIME_FUNCTION static
71  hipError_t SortPairs(void * d_temp_storage,
72  size_t& temp_storage_bytes,
73  DoubleBuffer<KeyT>& d_keys,
74  DoubleBuffer<ValueT>& d_values,
75  int num_items,
76  int num_segments,
77  OffsetIteratorT d_begin_offsets,
78  OffsetIteratorT d_end_offsets,
79  int begin_bit = 0,
80  int end_bit = sizeof(KeyT) * 8,
81  hipStream_t stream = 0,
82  bool debug_synchronous = false)
83  {
84  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
85  ::rocprim::double_buffer<ValueT> d_values_db = detail::to_double_buffer(d_values);
86  hipError_t error = ::rocprim::segmented_radix_sort_pairs(
87  d_temp_storage, temp_storage_bytes,
88  d_keys_db, d_values_db, num_items,
89  num_segments, d_begin_offsets, d_end_offsets,
90  begin_bit, end_bit,
91  stream, debug_synchronous
92  );
93  detail::update_double_buffer(d_keys, d_keys_db);
94  detail::update_double_buffer(d_values, d_values_db);
95  return error;
96  }
97 
98  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
99  HIPCUB_RUNTIME_FUNCTION static
100  hipError_t SortPairsDescending(void * d_temp_storage,
101  size_t& temp_storage_bytes,
102  const KeyT * d_keys_in,
103  KeyT * d_keys_out,
104  const ValueT * d_values_in,
105  ValueT * d_values_out,
106  int num_items,
107  int num_segments,
108  OffsetIteratorT d_begin_offsets,
109  OffsetIteratorT d_end_offsets,
110  int begin_bit = 0,
111  int end_bit = sizeof(KeyT) * 8,
112  hipStream_t stream = 0,
113  bool debug_synchronous = false)
114  {
115  return ::rocprim::segmented_radix_sort_pairs_desc(
116  d_temp_storage, temp_storage_bytes,
117  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,
118  num_segments, d_begin_offsets, d_end_offsets,
119  begin_bit, end_bit,
120  stream, debug_synchronous
121  );
122  }
123 
124  template<typename KeyT, typename ValueT, typename OffsetIteratorT>
125  HIPCUB_RUNTIME_FUNCTION static
126  hipError_t SortPairsDescending(void * d_temp_storage,
127  size_t& temp_storage_bytes,
128  DoubleBuffer<KeyT>& d_keys,
129  DoubleBuffer<ValueT>& d_values,
130  int num_items,
131  int num_segments,
132  OffsetIteratorT d_begin_offsets,
133  OffsetIteratorT d_end_offsets,
134  int begin_bit = 0,
135  int end_bit = sizeof(KeyT) * 8,
136  hipStream_t stream = 0,
137  bool debug_synchronous = false)
138  {
139  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
140  ::rocprim::double_buffer<ValueT> d_values_db = detail::to_double_buffer(d_values);
141  hipError_t error = ::rocprim::segmented_radix_sort_pairs_desc(
142  d_temp_storage, temp_storage_bytes,
143  d_keys_db, d_values_db, num_items,
144  num_segments, d_begin_offsets, d_end_offsets,
145  begin_bit, end_bit,
146  stream, debug_synchronous
147  );
148  detail::update_double_buffer(d_keys, d_keys_db);
149  detail::update_double_buffer(d_values, d_values_db);
150  return error;
151  }
152 
153  template<typename KeyT, typename OffsetIteratorT>
154  HIPCUB_RUNTIME_FUNCTION static
155  hipError_t SortKeys(void * d_temp_storage,
156  size_t& temp_storage_bytes,
157  const KeyT * d_keys_in,
158  KeyT * d_keys_out,
159  int num_items,
160  int num_segments,
161  OffsetIteratorT d_begin_offsets,
162  OffsetIteratorT d_end_offsets,
163  int begin_bit = 0,
164  int end_bit = sizeof(KeyT) * 8,
165  hipStream_t stream = 0,
166  bool debug_synchronous = false)
167  {
168  return ::rocprim::segmented_radix_sort_keys(
169  d_temp_storage, temp_storage_bytes,
170  d_keys_in, d_keys_out, num_items,
171  num_segments, d_begin_offsets, d_end_offsets,
172  begin_bit, end_bit,
173  stream, debug_synchronous
174  );
175  }
176 
177  template<typename KeyT, typename OffsetIteratorT>
178  HIPCUB_RUNTIME_FUNCTION static
179  hipError_t SortKeys(void * d_temp_storage,
180  size_t& temp_storage_bytes,
181  DoubleBuffer<KeyT>& d_keys,
182  int num_items,
183  int num_segments,
184  OffsetIteratorT d_begin_offsets,
185  OffsetIteratorT d_end_offsets,
186  int begin_bit = 0,
187  int end_bit = sizeof(KeyT) * 8,
188  hipStream_t stream = 0,
189  bool debug_synchronous = false)
190  {
191  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
192  hipError_t error = ::rocprim::segmented_radix_sort_keys(
193  d_temp_storage, temp_storage_bytes,
194  d_keys_db, num_items,
195  num_segments, d_begin_offsets, d_end_offsets,
196  begin_bit, end_bit,
197  stream, debug_synchronous
198  );
199  detail::update_double_buffer(d_keys, d_keys_db);
200  return error;
201  }
202 
203  template<typename KeyT, typename OffsetIteratorT>
204  HIPCUB_RUNTIME_FUNCTION static
205  hipError_t SortKeysDescending(void * d_temp_storage,
206  size_t& temp_storage_bytes,
207  const KeyT * d_keys_in,
208  KeyT * d_keys_out,
209  int num_items,
210  int num_segments,
211  OffsetIteratorT d_begin_offsets,
212  OffsetIteratorT d_end_offsets,
213  int begin_bit = 0,
214  int end_bit = sizeof(KeyT) * 8,
215  hipStream_t stream = 0,
216  bool debug_synchronous = false)
217  {
218  return ::rocprim::segmented_radix_sort_keys_desc(
219  d_temp_storage, temp_storage_bytes,
220  d_keys_in, d_keys_out, num_items,
221  num_segments, d_begin_offsets, d_end_offsets,
222  begin_bit, end_bit,
223  stream, debug_synchronous
224  );
225  }
226 
227  template<typename KeyT, typename OffsetIteratorT>
228  HIPCUB_RUNTIME_FUNCTION static
229  hipError_t SortKeysDescending(void * d_temp_storage,
230  size_t& temp_storage_bytes,
231  DoubleBuffer<KeyT>& d_keys,
232  int num_items,
233  int num_segments,
234  OffsetIteratorT d_begin_offsets,
235  OffsetIteratorT d_end_offsets,
236  int begin_bit = 0,
237  int end_bit = sizeof(KeyT) * 8,
238  hipStream_t stream = 0,
239  bool debug_synchronous = false)
240  {
241  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
242  hipError_t error = ::rocprim::segmented_radix_sort_keys_desc(
243  d_temp_storage, temp_storage_bytes,
244  d_keys_db, num_items,
245  num_segments, d_begin_offsets, d_end_offsets,
246  begin_bit, end_bit,
247  stream, debug_synchronous
248  );
249  detail::update_double_buffer(d_keys, d_keys_db);
250  return error;
251  }
252 };
253 
254 END_HIPCUB_NAMESPACE
255 
256 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_SEGMENTED_RADIX_SORT_HPP_