/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/device/device_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_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_radix_sort.hpp Source File
device_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_RADIX_SORT_HPP_
31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_RADIX_SORT_HPP_
32 
33 #include "../../../config.hpp"
34 
35 #include "../util_type.hpp"
36 
37 #include <rocprim/device/device_radix_sort.hpp>
38 
39 BEGIN_HIPCUB_NAMESPACE
40 
41 struct DeviceRadixSort
42 {
43  template<typename KeyT, typename ValueT, typename NumItemsT>
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  NumItemsT num_items,
52  int begin_bit = 0,
53  int end_bit = sizeof(KeyT) * 8,
54  hipStream_t stream = 0,
55  bool debug_synchronous = false)
56  {
57  return ::rocprim::radix_sort_pairs(
58  d_temp_storage, temp_storage_bytes,
59  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,
60  begin_bit, end_bit,
61  stream, debug_synchronous
62  );
63  }
64 
65  template<typename KeyT, typename ValueT, typename NumItemsT>
66  HIPCUB_RUNTIME_FUNCTION static
67  hipError_t SortPairs(void * d_temp_storage,
68  size_t& temp_storage_bytes,
69  DoubleBuffer<KeyT>& d_keys,
70  DoubleBuffer<ValueT>& d_values,
71  NumItemsT num_items,
72  int begin_bit = 0,
73  int end_bit = sizeof(KeyT) * 8,
74  hipStream_t stream = 0,
75  bool debug_synchronous = false)
76  {
77  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
78  ::rocprim::double_buffer<ValueT> d_values_db = detail::to_double_buffer(d_values);
79  hipError_t error = ::rocprim::radix_sort_pairs(
80  d_temp_storage, temp_storage_bytes,
81  d_keys_db, d_values_db, num_items,
82  begin_bit, end_bit,
83  stream, debug_synchronous
84  );
85  detail::update_double_buffer(d_keys, d_keys_db);
86  detail::update_double_buffer(d_values, d_values_db);
87  return error;
88  }
89 
90  template<typename KeyT, typename ValueT, typename NumItemsT>
91  HIPCUB_RUNTIME_FUNCTION static
92  hipError_t SortPairsDescending(void * d_temp_storage,
93  size_t& temp_storage_bytes,
94  const KeyT * d_keys_in,
95  KeyT * d_keys_out,
96  const ValueT * d_values_in,
97  ValueT * d_values_out,
98  NumItemsT num_items,
99  int begin_bit = 0,
100  int end_bit = sizeof(KeyT) * 8,
101  hipStream_t stream = 0,
102  bool debug_synchronous = false)
103  {
104  return ::rocprim::radix_sort_pairs_desc(
105  d_temp_storage, temp_storage_bytes,
106  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,
107  begin_bit, end_bit,
108  stream, debug_synchronous
109  );
110  }
111 
112  template<typename KeyT, typename ValueT, typename NumItemsT>
113  HIPCUB_RUNTIME_FUNCTION static
114  hipError_t SortPairsDescending(void * d_temp_storage,
115  size_t& temp_storage_bytes,
116  DoubleBuffer<KeyT>& d_keys,
117  DoubleBuffer<ValueT>& d_values,
118  NumItemsT num_items,
119  int begin_bit = 0,
120  int end_bit = sizeof(KeyT) * 8,
121  hipStream_t stream = 0,
122  bool debug_synchronous = false)
123  {
124  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
125  ::rocprim::double_buffer<ValueT> d_values_db = detail::to_double_buffer(d_values);
126  hipError_t error = ::rocprim::radix_sort_pairs_desc(
127  d_temp_storage, temp_storage_bytes,
128  d_keys_db, d_values_db, num_items,
129  begin_bit, end_bit,
130  stream, debug_synchronous
131  );
132  detail::update_double_buffer(d_keys, d_keys_db);
133  detail::update_double_buffer(d_values, d_values_db);
134  return error;
135  }
136 
137  template<typename KeyT, typename NumItemsT>
138  HIPCUB_RUNTIME_FUNCTION static
139  hipError_t SortKeys(void * d_temp_storage,
140  size_t& temp_storage_bytes,
141  const KeyT * d_keys_in,
142  KeyT * d_keys_out,
143  NumItemsT num_items,
144  int begin_bit = 0,
145  int end_bit = sizeof(KeyT) * 8,
146  hipStream_t stream = 0,
147  bool debug_synchronous = false)
148  {
149  return ::rocprim::radix_sort_keys(
150  d_temp_storage, temp_storage_bytes,
151  d_keys_in, d_keys_out, num_items,
152  begin_bit, end_bit,
153  stream, debug_synchronous
154  );
155  }
156 
157  template<typename KeyT, typename NumItemsT>
158  HIPCUB_RUNTIME_FUNCTION static
159  hipError_t SortKeys(void * d_temp_storage,
160  size_t& temp_storage_bytes,
161  DoubleBuffer<KeyT>& d_keys,
162  NumItemsT num_items,
163  int begin_bit = 0,
164  int end_bit = sizeof(KeyT) * 8,
165  hipStream_t stream = 0,
166  bool debug_synchronous = false)
167  {
168  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
169  hipError_t error = ::rocprim::radix_sort_keys(
170  d_temp_storage, temp_storage_bytes,
171  d_keys_db, num_items,
172  begin_bit, end_bit,
173  stream, debug_synchronous
174  );
175  detail::update_double_buffer(d_keys, d_keys_db);
176  return error;
177  }
178 
179  template<typename KeyT, typename NumItemsT>
180  HIPCUB_RUNTIME_FUNCTION static
181  hipError_t SortKeysDescending(void * d_temp_storage,
182  size_t& temp_storage_bytes,
183  const KeyT * d_keys_in,
184  KeyT * d_keys_out,
185  NumItemsT num_items,
186  int begin_bit = 0,
187  int end_bit = sizeof(KeyT) * 8,
188  hipStream_t stream = 0,
189  bool debug_synchronous = false)
190  {
191  return ::rocprim::radix_sort_keys_desc(
192  d_temp_storage, temp_storage_bytes,
193  d_keys_in, d_keys_out, num_items,
194  begin_bit, end_bit,
195  stream, debug_synchronous
196  );
197  }
198 
199  template<typename KeyT, typename NumItemsT>
200  HIPCUB_RUNTIME_FUNCTION static
201  hipError_t SortKeysDescending(void * d_temp_storage,
202  size_t& temp_storage_bytes,
203  DoubleBuffer<KeyT>& d_keys,
204  NumItemsT num_items,
205  int begin_bit = 0,
206  int end_bit = sizeof(KeyT) * 8,
207  hipStream_t stream = 0,
208  bool debug_synchronous = false)
209  {
210  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);
211  hipError_t error = ::rocprim::radix_sort_keys_desc(
212  d_temp_storage, temp_storage_bytes,
213  d_keys_db, num_items,
214  begin_bit, end_bit,
215  stream, debug_synchronous
216  );
217  detail::update_double_buffer(d_keys, d_keys_db);
218  return error;
219  }
220 };
221 
222 END_HIPCUB_NAMESPACE
223 
224 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_RADIX_SORT_HPP_