rocPRIM
device_radix_sort_onesweep.hpp
1 // Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved.
2 //
3 // Permission is hereby granted, free of charge, to any person obtaining a copy
4 // of this software and associated documentation files (the "Software"), to deal
5 // in the Software without restriction, including without limitation the rights
6 // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7 // copies of the Software, and to permit persons to whom the Software is
8 // furnished to do so, subject to the following conditions:
9 //
10 // The above copyright notice and this permission notice shall be included in
11 // all copies or substantial portions of the Software.
12 //
13 // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
14 // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
15 // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
16 // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
17 // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
18 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
19 // THE SOFTWARE.
20 
21 #ifndef ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_RADIX_SORT_ONESWEEP_HPP_
22 #define ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_RADIX_SORT_ONESWEEP_HPP_
23 
24 #include "../../../type_traits.hpp"
25 #include "../device_config_helper.hpp"
26 #include <type_traits>
27 
28 /* DO NOT EDIT THIS FILE
29  * This file is automatically generated by `/scripts/autotune/create_optimization.py`.
30  * so most likely you want to edit rocprim/device/device_(algo)_config.hpp
31  */
32 
35 
36 BEGIN_ROCPRIM_NAMESPACE
37 
38 namespace detail
39 {
40 
41 template<unsigned int arch,
42  class key_type,
43  class value_type = rocprim::empty_type,
44  class enable = void>
45 struct default_radix_sort_onesweep_config
46  : radix_sort_onesweep_config_base<key_type, value_type>::type
47 {};
48 
49 // Based on key_type = double, value_type = int64_t
50 template<class key_type, class value_type>
51 struct default_radix_sort_onesweep_config<
52  static_cast<unsigned int>(target_arch::gfx1030),
53  key_type,
54  value_type,
55  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
56  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
57  && (sizeof(value_type) > 4))>>
58  : radix_sort_onesweep_config<kernel_config<1024, 1>,
59  kernel_config<1024, 1>,
60  8,
61  block_radix_rank_algorithm::match>
62 {};
63 
64 // Based on key_type = double, value_type = int
65 template<class key_type, class value_type>
66 struct default_radix_sort_onesweep_config<
67  static_cast<unsigned int>(target_arch::gfx1030),
68  key_type,
69  value_type,
70  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
71  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
72  && (sizeof(value_type) > 2))>>
73  : radix_sort_onesweep_config<kernel_config<1024, 4>,
74  kernel_config<1024, 4>,
75  8,
76  block_radix_rank_algorithm::match>
77 {};
78 
79 // Based on key_type = double, value_type = short
80 template<class key_type, class value_type>
81 struct default_radix_sort_onesweep_config<
82  static_cast<unsigned int>(target_arch::gfx1030),
83  key_type,
84  value_type,
85  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
86  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
87  && (sizeof(value_type) > 1))>>
88  : radix_sort_onesweep_config<kernel_config<1024, 4>,
89  kernel_config<1024, 4>,
90  8,
91  block_radix_rank_algorithm::match>
92 {};
93 
94 // Based on key_type = double, value_type = int8_t
95 template<class key_type, class value_type>
96 struct default_radix_sort_onesweep_config<
97  static_cast<unsigned int>(target_arch::gfx1030),
98  key_type,
99  value_type,
100  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
101  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
102  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
103  : radix_sort_onesweep_config<kernel_config<1024, 4>,
104  kernel_config<1024, 4>,
105  8,
106  block_radix_rank_algorithm::match>
107 {};
108 
109 // Based on key_type = double, value_type = empty_type
110 template<class key_type, class value_type>
111 struct default_radix_sort_onesweep_config<
112  static_cast<unsigned int>(target_arch::gfx1030),
113  key_type,
114  value_type,
115  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
116  && (sizeof(key_type) > 4)
117  && (std::is_same<value_type, rocprim::empty_type>::value))>>
118  : radix_sort_onesweep_config<kernel_config<1024, 4>,
119  kernel_config<1024, 4>,
120  8,
121  block_radix_rank_algorithm::match>
122 {};
123 
124 // Based on key_type = float, value_type = int64_t
125 template<class key_type, class value_type>
126 struct default_radix_sort_onesweep_config<
127  static_cast<unsigned int>(target_arch::gfx1030),
128  key_type,
129  value_type,
130  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
131  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
132  && (sizeof(value_type) > 4))>>
133  : radix_sort_onesweep_config<kernel_config<1024, 4>,
134  kernel_config<1024, 4>,
135  8,
136  block_radix_rank_algorithm::match>
137 {};
138 
139 // Based on key_type = float, value_type = int
140 template<class key_type, class value_type>
141 struct default_radix_sort_onesweep_config<
142  static_cast<unsigned int>(target_arch::gfx1030),
143  key_type,
144  value_type,
145  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
146  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
147  && (sizeof(value_type) > 2))>>
148  : radix_sort_onesweep_config<kernel_config<1024, 8>,
149  kernel_config<1024, 8>,
150  8,
151  block_radix_rank_algorithm::match>
152 {};
153 
154 // Based on key_type = float, value_type = short
155 template<class key_type, class value_type>
156 struct default_radix_sort_onesweep_config<
157  static_cast<unsigned int>(target_arch::gfx1030),
158  key_type,
159  value_type,
160  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
161  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
162  && (sizeof(value_type) > 1))>>
163  : radix_sort_onesweep_config<kernel_config<1024, 4>,
164  kernel_config<1024, 4>,
165  8,
166  block_radix_rank_algorithm::match>
167 {};
168 
169 // Based on key_type = float, value_type = int8_t
170 template<class key_type, class value_type>
171 struct default_radix_sort_onesweep_config<
172  static_cast<unsigned int>(target_arch::gfx1030),
173  key_type,
174  value_type,
175  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
176  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
177  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
178  : radix_sort_onesweep_config<kernel_config<512, 16>,
179  kernel_config<512, 16>,
180  8,
181  block_radix_rank_algorithm::match>
182 {};
183 
184 // Based on key_type = float, value_type = empty_type
185 template<class key_type, class value_type>
186 struct default_radix_sort_onesweep_config<
187  static_cast<unsigned int>(target_arch::gfx1030),
188  key_type,
189  value_type,
190  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
191  && (sizeof(key_type) > 2)
192  && (std::is_same<value_type, rocprim::empty_type>::value))>>
193  : radix_sort_onesweep_config<kernel_config<1024, 6>,
194  kernel_config<1024, 6>,
195  8,
196  block_radix_rank_algorithm::match>
197 {};
198 
199 // Based on key_type = rocprim::half, value_type = int64_t
200 template<class key_type, class value_type>
201 struct default_radix_sort_onesweep_config<
202  static_cast<unsigned int>(target_arch::gfx1030),
203  key_type,
204  value_type,
205  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
206  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
207  : radix_sort_onesweep_config<kernel_config<512, 12>,
208  kernel_config<512, 12>,
209  8,
210  block_radix_rank_algorithm::match>
211 {};
212 
213 // Based on key_type = rocprim::half, value_type = int
214 template<class key_type, class value_type>
215 struct default_radix_sort_onesweep_config<
216  static_cast<unsigned int>(target_arch::gfx1030),
217  key_type,
218  value_type,
219  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
220  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
221  : radix_sort_onesweep_config<kernel_config<512, 16>,
222  kernel_config<512, 16>,
223  8,
224  block_radix_rank_algorithm::match>
225 {};
226 
227 // Based on key_type = rocprim::half, value_type = short
228 template<class key_type, class value_type>
229 struct default_radix_sort_onesweep_config<
230  static_cast<unsigned int>(target_arch::gfx1030),
231  key_type,
232  value_type,
233  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
234  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
235  : radix_sort_onesweep_config<kernel_config<512, 16>,
236  kernel_config<512, 16>,
237  8,
238  block_radix_rank_algorithm::match>
239 {};
240 
241 // Based on key_type = rocprim::half, value_type = int8_t
242 template<class key_type, class value_type>
243 struct default_radix_sort_onesweep_config<
244  static_cast<unsigned int>(target_arch::gfx1030),
245  key_type,
246  value_type,
247  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
248  && (sizeof(value_type) <= 1)
249  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
250  : radix_sort_onesweep_config<kernel_config<256, 18>,
251  kernel_config<256, 18>,
252  8,
253  block_radix_rank_algorithm::match>
254 {};
255 
256 // Based on key_type = rocprim::half, value_type = empty_type
257 template<class key_type, class value_type>
258 struct default_radix_sort_onesweep_config<
259  static_cast<unsigned int>(target_arch::gfx1030),
260  key_type,
261  value_type,
262  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
263  && (std::is_same<value_type, rocprim::empty_type>::value))>>
264  : radix_sort_onesweep_config<kernel_config<512, 12>,
265  kernel_config<512, 12>,
266  8,
267  block_radix_rank_algorithm::match>
268 {};
269 
270 // Based on key_type = int64_t, value_type = int64_t
271 template<class key_type, class value_type>
272 struct default_radix_sort_onesweep_config<
273  static_cast<unsigned int>(target_arch::gfx1030),
274  key_type,
275  value_type,
276  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
277  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
278  && (sizeof(value_type) > 4))>>
279  : radix_sort_onesweep_config<kernel_config<1024, 1>,
280  kernel_config<1024, 1>,
281  8,
282  block_radix_rank_algorithm::match>
283 {};
284 
285 // Based on key_type = int64_t, value_type = int
286 template<class key_type, class value_type>
287 struct default_radix_sort_onesweep_config<
288  static_cast<unsigned int>(target_arch::gfx1030),
289  key_type,
290  value_type,
291  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
292  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
293  && (sizeof(value_type) > 2))>>
294  : radix_sort_onesweep_config<kernel_config<1024, 1>,
295  kernel_config<1024, 1>,
296  8,
297  block_radix_rank_algorithm::match>
298 {};
299 
300 // Based on key_type = int64_t, value_type = short
301 template<class key_type, class value_type>
302 struct default_radix_sort_onesweep_config<
303  static_cast<unsigned int>(target_arch::gfx1030),
304  key_type,
305  value_type,
306  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
307  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
308  && (sizeof(value_type) > 1))>>
309  : radix_sort_onesweep_config<kernel_config<1024, 4>,
310  kernel_config<1024, 4>,
311  8,
312  block_radix_rank_algorithm::match>
313 {};
314 
315 // Based on key_type = int64_t, value_type = int8_t
316 template<class key_type, class value_type>
317 struct default_radix_sort_onesweep_config<
318  static_cast<unsigned int>(target_arch::gfx1030),
319  key_type,
320  value_type,
321  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
322  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
323  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
324  : radix_sort_onesweep_config<kernel_config<1024, 4>,
325  kernel_config<1024, 4>,
326  8,
327  block_radix_rank_algorithm::match>
328 {};
329 
330 // Based on key_type = int64_t, value_type = empty_type
331 template<class key_type, class value_type>
332 struct default_radix_sort_onesweep_config<
333  static_cast<unsigned int>(target_arch::gfx1030),
334  key_type,
335  value_type,
336  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
337  && (sizeof(key_type) > 4)
338  && (std::is_same<value_type, rocprim::empty_type>::value))>>
339  : radix_sort_onesweep_config<kernel_config<1024, 4>,
340  kernel_config<1024, 4>,
341  8,
342  block_radix_rank_algorithm::match>
343 {};
344 
345 // Based on key_type = int, value_type = int64_t
346 template<class key_type, class value_type>
347 struct default_radix_sort_onesweep_config<
348  static_cast<unsigned int>(target_arch::gfx1030),
349  key_type,
350  value_type,
351  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
352  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
353  && (sizeof(value_type) > 4))>>
354  : radix_sort_onesweep_config<kernel_config<1024, 1>,
355  kernel_config<1024, 1>,
356  8,
357  block_radix_rank_algorithm::match>
358 {};
359 
360 // Based on key_type = int, value_type = int
361 template<class key_type, class value_type>
362 struct default_radix_sort_onesweep_config<
363  static_cast<unsigned int>(target_arch::gfx1030),
364  key_type,
365  value_type,
366  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
367  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
368  && (sizeof(value_type) > 2))>>
369  : radix_sort_onesweep_config<kernel_config<1024, 4>,
370  kernel_config<1024, 4>,
371  7,
372  block_radix_rank_algorithm::match>
373 {};
374 
375 // Based on key_type = int, value_type = short
376 template<class key_type, class value_type>
377 struct default_radix_sort_onesweep_config<
378  static_cast<unsigned int>(target_arch::gfx1030),
379  key_type,
380  value_type,
381  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
382  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
383  && (sizeof(value_type) > 1))>>
384  : radix_sort_onesweep_config<kernel_config<1024, 4>,
385  kernel_config<1024, 4>,
386  8,
387  block_radix_rank_algorithm::match>
388 {};
389 
390 // Based on key_type = int, value_type = int8_t
391 template<class key_type, class value_type>
392 struct default_radix_sort_onesweep_config<
393  static_cast<unsigned int>(target_arch::gfx1030),
394  key_type,
395  value_type,
396  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
397  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
398  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
399  : radix_sort_onesweep_config<kernel_config<1024, 4>,
400  kernel_config<1024, 4>,
401  8,
402  block_radix_rank_algorithm::match>
403 {};
404 
405 // Based on key_type = int, value_type = empty_type
406 template<class key_type, class value_type>
407 struct default_radix_sort_onesweep_config<
408  static_cast<unsigned int>(target_arch::gfx1030),
409  key_type,
410  value_type,
411  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
412  && (sizeof(key_type) > 2)
413  && (std::is_same<value_type, rocprim::empty_type>::value))>>
414  : radix_sort_onesweep_config<kernel_config<1024, 4>,
415  kernel_config<1024, 4>,
416  8,
417  block_radix_rank_algorithm::match>
418 {};
419 
420 // Based on key_type = short, value_type = int64_t
421 template<class key_type, class value_type>
422 struct default_radix_sort_onesweep_config<
423  static_cast<unsigned int>(target_arch::gfx1030),
424  key_type,
425  value_type,
426  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
427  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
428  && (sizeof(value_type) > 4))>>
429  : radix_sort_onesweep_config<kernel_config<1024, 4>,
430  kernel_config<1024, 4>,
431  8,
432  block_radix_rank_algorithm::match>
433 {};
434 
435 // Based on key_type = short, value_type = int
436 template<class key_type, class value_type>
437 struct default_radix_sort_onesweep_config<
438  static_cast<unsigned int>(target_arch::gfx1030),
439  key_type,
440  value_type,
441  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
442  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
443  && (sizeof(value_type) > 2))>>
444  : radix_sort_onesweep_config<kernel_config<1024, 8>,
445  kernel_config<1024, 8>,
446  8,
447  block_radix_rank_algorithm::match>
448 {};
449 
450 // Based on key_type = short, value_type = short
451 template<class key_type, class value_type>
452 struct default_radix_sort_onesweep_config<
453  static_cast<unsigned int>(target_arch::gfx1030),
454  key_type,
455  value_type,
456  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
457  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
458  && (sizeof(value_type) > 1))>>
459  : radix_sort_onesweep_config<kernel_config<1024, 4>,
460  kernel_config<1024, 4>,
461  6,
462  block_radix_rank_algorithm::match>
463 {};
464 
465 // Based on key_type = short, value_type = int8_t
466 template<class key_type, class value_type>
467 struct default_radix_sort_onesweep_config<
468  static_cast<unsigned int>(target_arch::gfx1030),
469  key_type,
470  value_type,
471  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
472  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
473  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
474  : radix_sort_onesweep_config<kernel_config<512, 16>,
475  kernel_config<512, 16>,
476  8,
477  block_radix_rank_algorithm::match>
478 {};
479 
480 // Based on key_type = short, value_type = empty_type
481 template<class key_type, class value_type>
482 struct default_radix_sort_onesweep_config<
483  static_cast<unsigned int>(target_arch::gfx1030),
484  key_type,
485  value_type,
486  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
487  && (sizeof(key_type) > 1)
488  && (std::is_same<value_type, rocprim::empty_type>::value))>>
489  : radix_sort_onesweep_config<kernel_config<512, 22>,
490  kernel_config<512, 22>,
491  8,
492  block_radix_rank_algorithm::match>
493 {};
494 
495 // Based on key_type = int8_t, value_type = int64_t
496 template<class key_type, class value_type>
497 struct default_radix_sort_onesweep_config<
498  static_cast<unsigned int>(target_arch::gfx1030),
499  key_type,
500  value_type,
501  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
502  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
503  : radix_sort_onesweep_config<kernel_config<512, 12>,
504  kernel_config<512, 12>,
505  8,
506  block_radix_rank_algorithm::match>
507 {};
508 
509 // Based on key_type = int8_t, value_type = int
510 template<class key_type, class value_type>
511 struct default_radix_sort_onesweep_config<
512  static_cast<unsigned int>(target_arch::gfx1030),
513  key_type,
514  value_type,
515  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
516  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
517  : radix_sort_onesweep_config<kernel_config<1024, 12>,
518  kernel_config<1024, 12>,
519  8,
520  block_radix_rank_algorithm::match>
521 {};
522 
523 // Based on key_type = int8_t, value_type = short
524 template<class key_type, class value_type>
525 struct default_radix_sort_onesweep_config<
526  static_cast<unsigned int>(target_arch::gfx1030),
527  key_type,
528  value_type,
529  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
530  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
531  : radix_sort_onesweep_config<kernel_config<1024, 12>,
532  kernel_config<1024, 12>,
533  8,
534  block_radix_rank_algorithm::match>
535 {};
536 
537 // Based on key_type = int8_t, value_type = int8_t
538 template<class key_type, class value_type>
539 struct default_radix_sort_onesweep_config<
540  static_cast<unsigned int>(target_arch::gfx1030),
541  key_type,
542  value_type,
543  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
544  && (sizeof(value_type) <= 1)
545  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
546  : radix_sort_onesweep_config<kernel_config<256, 22>,
547  kernel_config<256, 22>,
548  8,
549  block_radix_rank_algorithm::match>
550 {};
551 
552 // Based on key_type = int8_t, value_type = empty_type
553 template<class key_type, class value_type>
554 struct default_radix_sort_onesweep_config<
555  static_cast<unsigned int>(target_arch::gfx1030),
556  key_type,
557  value_type,
558  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
559  && (std::is_same<value_type, rocprim::empty_type>::value))>>
560  : radix_sort_onesweep_config<kernel_config<256, 18>,
561  kernel_config<256, 18>,
562  8,
563  block_radix_rank_algorithm::match>
564 {};
565 
566 // Based on key_type = double, value_type = int64_t
567 template<class key_type, class value_type>
568 struct default_radix_sort_onesweep_config<
569  static_cast<unsigned int>(target_arch::gfx900),
570  key_type,
571  value_type,
572  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
573  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
574  && (sizeof(value_type) > 4))>>
575  : radix_sort_onesweep_config<kernel_config<512, 8>,
576  kernel_config<512, 8>,
577  8,
578  block_radix_rank_algorithm::match>
579 {};
580 
581 // Based on key_type = double, value_type = int
582 template<class key_type, class value_type>
583 struct default_radix_sort_onesweep_config<
584  static_cast<unsigned int>(target_arch::gfx900),
585  key_type,
586  value_type,
587  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
588  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
589  && (sizeof(value_type) > 2))>>
590  : radix_sort_onesweep_config<kernel_config<512, 4>,
591  kernel_config<512, 4>,
592  8,
593  block_radix_rank_algorithm::match>
594 {};
595 
596 // Based on key_type = double, value_type = short
597 template<class key_type, class value_type>
598 struct default_radix_sort_onesweep_config<
599  static_cast<unsigned int>(target_arch::gfx900),
600  key_type,
601  value_type,
602  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
603  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
604  && (sizeof(value_type) > 1))>>
605  : radix_sort_onesweep_config<kernel_config<512, 12>,
606  kernel_config<512, 12>,
607  8,
608  block_radix_rank_algorithm::match>
609 {};
610 
611 // Based on key_type = double, value_type = int8_t
612 template<class key_type, class value_type>
613 struct default_radix_sort_onesweep_config<
614  static_cast<unsigned int>(target_arch::gfx900),
615  key_type,
616  value_type,
617  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
618  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
619  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
620  : radix_sort_onesweep_config<kernel_config<512, 6>,
621  kernel_config<512, 6>,
622  8,
623  block_radix_rank_algorithm::match>
624 {};
625 
626 // Based on key_type = double, value_type = empty_type
627 template<class key_type, class value_type>
628 struct default_radix_sort_onesweep_config<
629  static_cast<unsigned int>(target_arch::gfx900),
630  key_type,
631  value_type,
632  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
633  && (sizeof(key_type) > 4)
634  && (std::is_same<value_type, rocprim::empty_type>::value))>>
635  : radix_sort_onesweep_config<kernel_config<512, 12>,
636  kernel_config<512, 12>,
637  8,
638  block_radix_rank_algorithm::match>
639 {};
640 
641 // Based on key_type = float, value_type = int64_t
642 template<class key_type, class value_type>
643 struct default_radix_sort_onesweep_config<
644  static_cast<unsigned int>(target_arch::gfx900),
645  key_type,
646  value_type,
647  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
648  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
649  && (sizeof(value_type) > 4))>>
650  : radix_sort_onesweep_config<kernel_config<512, 8>,
651  kernel_config<512, 8>,
652  8,
653  block_radix_rank_algorithm::match>
654 {};
655 
656 // Based on key_type = float, value_type = int
657 template<class key_type, class value_type>
658 struct default_radix_sort_onesweep_config<
659  static_cast<unsigned int>(target_arch::gfx900),
660  key_type,
661  value_type,
662  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
663  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
664  && (sizeof(value_type) > 2))>>
665  : radix_sort_onesweep_config<kernel_config<512, 16>,
666  kernel_config<512, 16>,
667  8,
668  block_radix_rank_algorithm::match>
669 {};
670 
671 // Based on key_type = float, value_type = short
672 template<class key_type, class value_type>
673 struct default_radix_sort_onesweep_config<
674  static_cast<unsigned int>(target_arch::gfx900),
675  key_type,
676  value_type,
677  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
678  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
679  && (sizeof(value_type) > 1))>>
680  : radix_sort_onesweep_config<kernel_config<512, 6>,
681  kernel_config<512, 6>,
682  8,
683  block_radix_rank_algorithm::match>
684 {};
685 
686 // Based on key_type = float, value_type = int8_t
687 template<class key_type, class value_type>
688 struct default_radix_sort_onesweep_config<
689  static_cast<unsigned int>(target_arch::gfx900),
690  key_type,
691  value_type,
692  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
693  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
694  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
695  : radix_sort_onesweep_config<kernel_config<512, 8>,
696  kernel_config<512, 8>,
697  8,
698  block_radix_rank_algorithm::match>
699 {};
700 
701 // Based on key_type = float, value_type = empty_type
702 template<class key_type, class value_type>
703 struct default_radix_sort_onesweep_config<
704  static_cast<unsigned int>(target_arch::gfx900),
705  key_type,
706  value_type,
707  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
708  && (sizeof(key_type) > 2)
709  && (std::is_same<value_type, rocprim::empty_type>::value))>>
710  : radix_sort_onesweep_config<kernel_config<512, 8>,
711  kernel_config<512, 8>,
712  8,
713  block_radix_rank_algorithm::match>
714 {};
715 
716 // Based on key_type = rocprim::half, value_type = int64_t
717 template<class key_type, class value_type>
718 struct default_radix_sort_onesweep_config<
719  static_cast<unsigned int>(target_arch::gfx900),
720  key_type,
721  value_type,
722  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
723  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
724  : radix_sort_onesweep_config<kernel_config<512, 6>,
725  kernel_config<512, 6>,
726  8,
727  block_radix_rank_algorithm::match>
728 {};
729 
730 // Based on key_type = rocprim::half, value_type = int
731 template<class key_type, class value_type>
732 struct default_radix_sort_onesweep_config<
733  static_cast<unsigned int>(target_arch::gfx900),
734  key_type,
735  value_type,
736  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
737  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
738  : radix_sort_onesweep_config<kernel_config<256, 16>,
739  kernel_config<256, 16>,
740  8,
741  block_radix_rank_algorithm::match>
742 {};
743 
744 // Based on key_type = rocprim::half, value_type = short
745 template<class key_type, class value_type>
746 struct default_radix_sort_onesweep_config<
747  static_cast<unsigned int>(target_arch::gfx900),
748  key_type,
749  value_type,
750  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
751  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
752  : radix_sort_onesweep_config<kernel_config<256, 16>,
753  kernel_config<256, 16>,
754  8,
755  block_radix_rank_algorithm::match>
756 {};
757 
758 // Based on key_type = rocprim::half, value_type = int8_t
759 template<class key_type, class value_type>
760 struct default_radix_sort_onesweep_config<
761  static_cast<unsigned int>(target_arch::gfx900),
762  key_type,
763  value_type,
764  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
765  && (sizeof(value_type) <= 1)
766  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
767  : radix_sort_onesweep_config<kernel_config<256, 16>,
768  kernel_config<256, 16>,
769  8,
770  block_radix_rank_algorithm::match>
771 {};
772 
773 // Based on key_type = rocprim::half, value_type = empty_type
774 template<class key_type, class value_type>
775 struct default_radix_sort_onesweep_config<
776  static_cast<unsigned int>(target_arch::gfx900),
777  key_type,
778  value_type,
779  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
780  && (std::is_same<value_type, rocprim::empty_type>::value))>>
781  : radix_sort_onesweep_config<kernel_config<256, 16>,
782  kernel_config<256, 16>,
783  8,
784  block_radix_rank_algorithm::match>
785 {};
786 
787 // Based on key_type = int64_t, value_type = int64_t
788 template<class key_type, class value_type>
789 struct default_radix_sort_onesweep_config<
790  static_cast<unsigned int>(target_arch::gfx900),
791  key_type,
792  value_type,
793  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
794  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
795  && (sizeof(value_type) > 4))>>
796  : radix_sort_onesweep_config<kernel_config<256, 22>,
797  kernel_config<256, 22>,
798  8,
799  block_radix_rank_algorithm::match>
800 {};
801 
802 // Based on key_type = int64_t, value_type = int
803 template<class key_type, class value_type>
804 struct default_radix_sort_onesweep_config<
805  static_cast<unsigned int>(target_arch::gfx900),
806  key_type,
807  value_type,
808  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
809  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
810  && (sizeof(value_type) > 2))>>
811  : radix_sort_onesweep_config<kernel_config<512, 8>,
812  kernel_config<512, 8>,
813  8,
814  block_radix_rank_algorithm::match>
815 {};
816 
817 // Based on key_type = int64_t, value_type = short
818 template<class key_type, class value_type>
819 struct default_radix_sort_onesweep_config<
820  static_cast<unsigned int>(target_arch::gfx900),
821  key_type,
822  value_type,
823  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
824  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
825  && (sizeof(value_type) > 1))>>
826  : radix_sort_onesweep_config<kernel_config<512, 12>,
827  kernel_config<512, 12>,
828  8,
829  block_radix_rank_algorithm::match>
830 {};
831 
832 // Based on key_type = int64_t, value_type = int8_t
833 template<class key_type, class value_type>
834 struct default_radix_sort_onesweep_config<
835  static_cast<unsigned int>(target_arch::gfx900),
836  key_type,
837  value_type,
838  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
839  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
840  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
841  : radix_sort_onesweep_config<kernel_config<256, 18>,
842  kernel_config<256, 18>,
843  8,
844  block_radix_rank_algorithm::match>
845 {};
846 
847 // Based on key_type = int64_t, value_type = empty_type
848 template<class key_type, class value_type>
849 struct default_radix_sort_onesweep_config<
850  static_cast<unsigned int>(target_arch::gfx900),
851  key_type,
852  value_type,
853  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
854  && (sizeof(key_type) > 4)
855  && (std::is_same<value_type, rocprim::empty_type>::value))>>
856  : radix_sort_onesweep_config<kernel_config<1024, 4>,
857  kernel_config<1024, 4>,
858  8,
859  block_radix_rank_algorithm::match>
860 {};
861 
862 // Based on key_type = int, value_type = int64_t
863 template<class key_type, class value_type>
864 struct default_radix_sort_onesweep_config<
865  static_cast<unsigned int>(target_arch::gfx900),
866  key_type,
867  value_type,
868  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
869  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
870  && (sizeof(value_type) > 4))>>
871  : radix_sort_onesweep_config<kernel_config<256, 16>,
872  kernel_config<256, 16>,
873  8,
874  block_radix_rank_algorithm::match>
875 {};
876 
877 // Based on key_type = int, value_type = int
878 template<class key_type, class value_type>
879 struct default_radix_sort_onesweep_config<
880  static_cast<unsigned int>(target_arch::gfx900),
881  key_type,
882  value_type,
883  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
884  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
885  && (sizeof(value_type) > 2))>>
886  : radix_sort_onesweep_config<kernel_config<1024, 12>,
887  kernel_config<1024, 12>,
888  4,
889  block_radix_rank_algorithm::basic>
890 {};
891 
892 // Based on key_type = int, value_type = short
893 template<class key_type, class value_type>
894 struct default_radix_sort_onesweep_config<
895  static_cast<unsigned int>(target_arch::gfx900),
896  key_type,
897  value_type,
898  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
899  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
900  && (sizeof(value_type) > 1))>>
901  : radix_sort_onesweep_config<kernel_config<512, 16>,
902  kernel_config<512, 16>,
903  8,
904  block_radix_rank_algorithm::match>
905 {};
906 
907 // Based on key_type = int, value_type = int8_t
908 template<class key_type, class value_type>
909 struct default_radix_sort_onesweep_config<
910  static_cast<unsigned int>(target_arch::gfx900),
911  key_type,
912  value_type,
913  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
914  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
915  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
916  : radix_sort_onesweep_config<kernel_config<512, 22>,
917  kernel_config<512, 22>,
918  8,
919  block_radix_rank_algorithm::match>
920 {};
921 
922 // Based on key_type = int, value_type = empty_type
923 template<class key_type, class value_type>
924 struct default_radix_sort_onesweep_config<
925  static_cast<unsigned int>(target_arch::gfx900),
926  key_type,
927  value_type,
928  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
929  && (sizeof(key_type) > 2)
930  && (std::is_same<value_type, rocprim::empty_type>::value))>>
931  : radix_sort_onesweep_config<kernel_config<256, 18>,
932  kernel_config<256, 18>,
933  8,
934  block_radix_rank_algorithm::match>
935 {};
936 
937 // Based on key_type = short, value_type = int64_t
938 template<class key_type, class value_type>
939 struct default_radix_sort_onesweep_config<
940  static_cast<unsigned int>(target_arch::gfx900),
941  key_type,
942  value_type,
943  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
944  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
945  && (sizeof(value_type) > 4))>>
946  : radix_sort_onesweep_config<kernel_config<256, 12>,
947  kernel_config<256, 12>,
948  8,
949  block_radix_rank_algorithm::match>
950 {};
951 
952 // Based on key_type = short, value_type = int
953 template<class key_type, class value_type>
954 struct default_radix_sort_onesweep_config<
955  static_cast<unsigned int>(target_arch::gfx900),
956  key_type,
957  value_type,
958  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
959  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
960  && (sizeof(value_type) > 2))>>
961  : radix_sort_onesweep_config<kernel_config<256, 8>,
962  kernel_config<256, 8>,
963  8,
964  block_radix_rank_algorithm::match>
965 {};
966 
967 // Based on key_type = short, value_type = short
968 template<class key_type, class value_type>
969 struct default_radix_sort_onesweep_config<
970  static_cast<unsigned int>(target_arch::gfx900),
971  key_type,
972  value_type,
973  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
974  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
975  && (sizeof(value_type) > 1))>>
976  : radix_sort_onesweep_config<kernel_config<256, 18>,
977  kernel_config<256, 18>,
978  8,
979  block_radix_rank_algorithm::match>
980 {};
981 
982 // Based on key_type = short, value_type = int8_t
983 template<class key_type, class value_type>
984 struct default_radix_sort_onesweep_config<
985  static_cast<unsigned int>(target_arch::gfx900),
986  key_type,
987  value_type,
988  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
989  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
990  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
991  : radix_sort_onesweep_config<kernel_config<512, 8>,
992  kernel_config<512, 8>,
993  8,
994  block_radix_rank_algorithm::match>
995 {};
996 
997 // Based on key_type = short, value_type = empty_type
998 template<class key_type, class value_type>
999 struct default_radix_sort_onesweep_config<
1000  static_cast<unsigned int>(target_arch::gfx900),
1001  key_type,
1002  value_type,
1003  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1004  && (sizeof(key_type) > 1)
1005  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1006  : radix_sort_onesweep_config<kernel_config<512, 12>,
1007  kernel_config<512, 12>,
1008  4,
1009  block_radix_rank_algorithm::match>
1010 {};
1011 
1012 // Based on key_type = int8_t, value_type = int64_t
1013 template<class key_type, class value_type>
1014 struct default_radix_sort_onesweep_config<
1015  static_cast<unsigned int>(target_arch::gfx900),
1016  key_type,
1017  value_type,
1018  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1019  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1020  : radix_sort_onesweep_config<kernel_config<256, 22>,
1021  kernel_config<256, 22>,
1022  8,
1023  block_radix_rank_algorithm::match>
1024 {};
1025 
1026 // Based on key_type = int8_t, value_type = int
1027 template<class key_type, class value_type>
1028 struct default_radix_sort_onesweep_config<
1029  static_cast<unsigned int>(target_arch::gfx900),
1030  key_type,
1031  value_type,
1032  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1033  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1034  : radix_sort_onesweep_config<kernel_config<512, 22>,
1035  kernel_config<512, 22>,
1036  8,
1037  block_radix_rank_algorithm::match>
1038 {};
1039 
1040 // Based on key_type = int8_t, value_type = short
1041 template<class key_type, class value_type>
1042 struct default_radix_sort_onesweep_config<
1043  static_cast<unsigned int>(target_arch::gfx900),
1044  key_type,
1045  value_type,
1046  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1047  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1048  : radix_sort_onesweep_config<kernel_config<256, 18>,
1049  kernel_config<256, 18>,
1050  8,
1051  block_radix_rank_algorithm::match>
1052 {};
1053 
1054 // Based on key_type = int8_t, value_type = int8_t
1055 template<class key_type, class value_type>
1056 struct default_radix_sort_onesweep_config<
1057  static_cast<unsigned int>(target_arch::gfx900),
1058  key_type,
1059  value_type,
1060  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1061  && (sizeof(value_type) <= 1)
1062  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1063  : radix_sort_onesweep_config<kernel_config<256, 22>,
1064  kernel_config<256, 22>,
1065  8,
1066  block_radix_rank_algorithm::match>
1067 {};
1068 
1069 // Based on key_type = int8_t, value_type = empty_type
1070 template<class key_type, class value_type>
1071 struct default_radix_sort_onesweep_config<
1072  static_cast<unsigned int>(target_arch::gfx900),
1073  key_type,
1074  value_type,
1075  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1076  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1077  : radix_sort_onesweep_config<kernel_config<256, 16>,
1078  kernel_config<256, 16>,
1079  8,
1080  block_radix_rank_algorithm::match>
1081 {};
1082 
1083 // Based on key_type = double, value_type = int64_t
1084 template<class key_type, class value_type>
1085 struct default_radix_sort_onesweep_config<
1086  static_cast<unsigned int>(target_arch::gfx908),
1087  key_type,
1088  value_type,
1089  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1090  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1091  && (sizeof(value_type) > 4))>>
1092  : radix_sort_onesweep_config<kernel_config<1024, 6>,
1093  kernel_config<1024, 6>,
1094  8,
1095  block_radix_rank_algorithm::match>
1096 {};
1097 
1098 // Based on key_type = double, value_type = int
1099 template<class key_type, class value_type>
1100 struct default_radix_sort_onesweep_config<
1101  static_cast<unsigned int>(target_arch::gfx908),
1102  key_type,
1103  value_type,
1104  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1105  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1106  && (sizeof(value_type) > 2))>>
1107  : radix_sort_onesweep_config<kernel_config<256, 12>,
1108  kernel_config<256, 12>,
1109  8,
1110  block_radix_rank_algorithm::match>
1111 {};
1112 
1113 // Based on key_type = double, value_type = short
1114 template<class key_type, class value_type>
1115 struct default_radix_sort_onesweep_config<
1116  static_cast<unsigned int>(target_arch::gfx908),
1117  key_type,
1118  value_type,
1119  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1120  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
1121  && (sizeof(value_type) > 1))>>
1122  : radix_sort_onesweep_config<kernel_config<256, 12>,
1123  kernel_config<256, 12>,
1124  8,
1125  block_radix_rank_algorithm::match>
1126 {};
1127 
1128 // Based on key_type = double, value_type = int8_t
1129 template<class key_type, class value_type>
1130 struct default_radix_sort_onesweep_config<
1131  static_cast<unsigned int>(target_arch::gfx908),
1132  key_type,
1133  value_type,
1134  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1135  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
1136  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1137  : radix_sort_onesweep_config<kernel_config<512, 6>,
1138  kernel_config<512, 6>,
1139  8,
1140  block_radix_rank_algorithm::match>
1141 {};
1142 
1143 // Based on key_type = double, value_type = empty_type
1144 template<class key_type, class value_type>
1145 struct default_radix_sort_onesweep_config<
1146  static_cast<unsigned int>(target_arch::gfx908),
1147  key_type,
1148  value_type,
1149  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1150  && (sizeof(key_type) > 4)
1151  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1152  : radix_sort_onesweep_config<kernel_config<512, 6>,
1153  kernel_config<512, 6>,
1154  8,
1155  block_radix_rank_algorithm::match>
1156 {};
1157 
1158 // Based on key_type = float, value_type = int64_t
1159 template<class key_type, class value_type>
1160 struct default_radix_sort_onesweep_config<
1161  static_cast<unsigned int>(target_arch::gfx908),
1162  key_type,
1163  value_type,
1164  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1165  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
1166  && (sizeof(value_type) > 4))>>
1167  : radix_sort_onesweep_config<kernel_config<1024, 6>,
1168  kernel_config<1024, 6>,
1169  8,
1170  block_radix_rank_algorithm::match>
1171 {};
1172 
1173 // Based on key_type = float, value_type = int
1174 template<class key_type, class value_type>
1175 struct default_radix_sort_onesweep_config<
1176  static_cast<unsigned int>(target_arch::gfx908),
1177  key_type,
1178  value_type,
1179  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1180  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1181  && (sizeof(value_type) > 2))>>
1182  : radix_sort_onesweep_config<kernel_config<256, 22>,
1183  kernel_config<256, 22>,
1184  8,
1185  block_radix_rank_algorithm::match>
1186 {};
1187 
1188 // Based on key_type = float, value_type = short
1189 template<class key_type, class value_type>
1190 struct default_radix_sort_onesweep_config<
1191  static_cast<unsigned int>(target_arch::gfx908),
1192  key_type,
1193  value_type,
1194  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1195  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1196  && (sizeof(value_type) > 1))>>
1197  : radix_sort_onesweep_config<kernel_config<256, 22>,
1198  kernel_config<256, 22>,
1199  8,
1200  block_radix_rank_algorithm::match>
1201 {};
1202 
1203 // Based on key_type = float, value_type = int8_t
1204 template<class key_type, class value_type>
1205 struct default_radix_sort_onesweep_config<
1206  static_cast<unsigned int>(target_arch::gfx908),
1207  key_type,
1208  value_type,
1209  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1210  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
1211  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1212  : radix_sort_onesweep_config<kernel_config<256, 22>,
1213  kernel_config<256, 22>,
1214  8,
1215  block_radix_rank_algorithm::match>
1216 {};
1217 
1218 // Based on key_type = float, value_type = empty_type
1219 template<class key_type, class value_type>
1220 struct default_radix_sort_onesweep_config<
1221  static_cast<unsigned int>(target_arch::gfx908),
1222  key_type,
1223  value_type,
1224  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1225  && (sizeof(key_type) > 2)
1226  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1227  : radix_sort_onesweep_config<kernel_config<256, 22>,
1228  kernel_config<256, 22>,
1229  8,
1230  block_radix_rank_algorithm::match>
1231 {};
1232 
1233 // Based on key_type = rocprim::half, value_type = int64_t
1234 template<class key_type, class value_type>
1235 struct default_radix_sort_onesweep_config<
1236  static_cast<unsigned int>(target_arch::gfx908),
1237  key_type,
1238  value_type,
1239  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1240  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1241  : radix_sort_onesweep_config<kernel_config<256, 12>,
1242  kernel_config<256, 12>,
1243  8,
1244  block_radix_rank_algorithm::match>
1245 {};
1246 
1247 // Based on key_type = rocprim::half, value_type = int
1248 template<class key_type, class value_type>
1249 struct default_radix_sort_onesweep_config<
1250  static_cast<unsigned int>(target_arch::gfx908),
1251  key_type,
1252  value_type,
1253  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1254  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1255  : radix_sort_onesweep_config<kernel_config<256, 16>,
1256  kernel_config<256, 16>,
1257  8,
1258  block_radix_rank_algorithm::match>
1259 {};
1260 
1261 // Based on key_type = rocprim::half, value_type = short
1262 template<class key_type, class value_type>
1263 struct default_radix_sort_onesweep_config<
1264  static_cast<unsigned int>(target_arch::gfx908),
1265  key_type,
1266  value_type,
1267  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1268  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1269  : radix_sort_onesweep_config<kernel_config<256, 16>,
1270  kernel_config<256, 16>,
1271  8,
1272  block_radix_rank_algorithm::match>
1273 {};
1274 
1275 // Based on key_type = rocprim::half, value_type = int8_t
1276 template<class key_type, class value_type>
1277 struct default_radix_sort_onesweep_config<
1278  static_cast<unsigned int>(target_arch::gfx908),
1279  key_type,
1280  value_type,
1281  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1282  && (sizeof(value_type) <= 1)
1283  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1284  : radix_sort_onesweep_config<kernel_config<256, 16>,
1285  kernel_config<256, 16>,
1286  8,
1287  block_radix_rank_algorithm::match>
1288 {};
1289 
1290 // Based on key_type = rocprim::half, value_type = empty_type
1291 template<class key_type, class value_type>
1292 struct default_radix_sort_onesweep_config<
1293  static_cast<unsigned int>(target_arch::gfx908),
1294  key_type,
1295  value_type,
1296  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1297  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1298  : radix_sort_onesweep_config<kernel_config<256, 16>,
1299  kernel_config<256, 16>,
1300  8,
1301  block_radix_rank_algorithm::match>
1302 {};
1303 
1304 // Based on key_type = int64_t, value_type = int64_t
1305 template<class key_type, class value_type>
1306 struct default_radix_sort_onesweep_config<
1307  static_cast<unsigned int>(target_arch::gfx908),
1308  key_type,
1309  value_type,
1310  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1311  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1312  && (sizeof(value_type) > 4))>>
1313  : radix_sort_onesweep_config<kernel_config<1024, 6>,
1314  kernel_config<1024, 6>,
1315  6,
1316  block_radix_rank_algorithm::match>
1317 {};
1318 
1319 // Based on key_type = int64_t, value_type = int
1320 template<class key_type, class value_type>
1321 struct default_radix_sort_onesweep_config<
1322  static_cast<unsigned int>(target_arch::gfx908),
1323  key_type,
1324  value_type,
1325  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1326  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1327  && (sizeof(value_type) > 2))>>
1328  : radix_sort_onesweep_config<kernel_config<256, 12>,
1329  kernel_config<256, 12>,
1330  8,
1331  block_radix_rank_algorithm::match>
1332 {};
1333 
1334 // Based on key_type = int64_t, value_type = short
1335 template<class key_type, class value_type>
1336 struct default_radix_sort_onesweep_config<
1337  static_cast<unsigned int>(target_arch::gfx908),
1338  key_type,
1339  value_type,
1340  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1341  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
1342  && (sizeof(value_type) > 1))>>
1343  : radix_sort_onesweep_config<kernel_config<256, 12>,
1344  kernel_config<256, 12>,
1345  8,
1346  block_radix_rank_algorithm::match>
1347 {};
1348 
1349 // Based on key_type = int64_t, value_type = int8_t
1350 template<class key_type, class value_type>
1351 struct default_radix_sort_onesweep_config<
1352  static_cast<unsigned int>(target_arch::gfx908),
1353  key_type,
1354  value_type,
1355  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1356  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
1357  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1358  : radix_sort_onesweep_config<kernel_config<512, 6>,
1359  kernel_config<512, 6>,
1360  8,
1361  block_radix_rank_algorithm::match>
1362 {};
1363 
1364 // Based on key_type = int64_t, value_type = empty_type
1365 template<class key_type, class value_type>
1366 struct default_radix_sort_onesweep_config<
1367  static_cast<unsigned int>(target_arch::gfx908),
1368  key_type,
1369  value_type,
1370  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1371  && (sizeof(key_type) > 4)
1372  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1373  : radix_sort_onesweep_config<kernel_config<256, 12>,
1374  kernel_config<256, 12>,
1375  8,
1376  block_radix_rank_algorithm::match>
1377 {};
1378 
1379 // Based on key_type = int, value_type = int64_t
1380 template<class key_type, class value_type>
1381 struct default_radix_sort_onesweep_config<
1382  static_cast<unsigned int>(target_arch::gfx908),
1383  key_type,
1384  value_type,
1385  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1386  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
1387  && (sizeof(value_type) > 4))>>
1388  : radix_sort_onesweep_config<kernel_config<256, 12>,
1389  kernel_config<256, 12>,
1390  8,
1391  block_radix_rank_algorithm::match>
1392 {};
1393 
1394 // Based on key_type = int, value_type = int
1395 template<class key_type, class value_type>
1396 struct default_radix_sort_onesweep_config<
1397  static_cast<unsigned int>(target_arch::gfx908),
1398  key_type,
1399  value_type,
1400  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1401  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1402  && (sizeof(value_type) > 2))>>
1403  : radix_sort_onesweep_config<kernel_config<512, 12>,
1404  kernel_config<512, 12>,
1405  5,
1406  block_radix_rank_algorithm::match>
1407 {};
1408 
1409 // Based on key_type = int, value_type = short
1410 template<class key_type, class value_type>
1411 struct default_radix_sort_onesweep_config<
1412  static_cast<unsigned int>(target_arch::gfx908),
1413  key_type,
1414  value_type,
1415  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1416  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1417  && (sizeof(value_type) > 1))>>
1418  : radix_sort_onesweep_config<kernel_config<256, 16>,
1419  kernel_config<256, 16>,
1420  8,
1421  block_radix_rank_algorithm::match>
1422 {};
1423 
1424 // Based on key_type = int, value_type = int8_t
1425 template<class key_type, class value_type>
1426 struct default_radix_sort_onesweep_config<
1427  static_cast<unsigned int>(target_arch::gfx908),
1428  key_type,
1429  value_type,
1430  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1431  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
1432  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1433  : radix_sort_onesweep_config<kernel_config<256, 22>,
1434  kernel_config<256, 22>,
1435  8,
1436  block_radix_rank_algorithm::match>
1437 {};
1438 
1439 // Based on key_type = int, value_type = empty_type
1440 template<class key_type, class value_type>
1441 struct default_radix_sort_onesweep_config<
1442  static_cast<unsigned int>(target_arch::gfx908),
1443  key_type,
1444  value_type,
1445  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1446  && (sizeof(key_type) > 2)
1447  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1448  : radix_sort_onesweep_config<kernel_config<256, 22>,
1449  kernel_config<256, 22>,
1450  7,
1451  block_radix_rank_algorithm::match>
1452 {};
1453 
1454 // Based on key_type = short, value_type = int64_t
1455 template<class key_type, class value_type>
1456 struct default_radix_sort_onesweep_config<
1457  static_cast<unsigned int>(target_arch::gfx908),
1458  key_type,
1459  value_type,
1460  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1461  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
1462  && (sizeof(value_type) > 4))>>
1463  : radix_sort_onesweep_config<kernel_config<256, 18>,
1464  kernel_config<256, 18>,
1465  8,
1466  block_radix_rank_algorithm::match>
1467 {};
1468 
1469 // Based on key_type = short, value_type = int
1470 template<class key_type, class value_type>
1471 struct default_radix_sort_onesweep_config<
1472  static_cast<unsigned int>(target_arch::gfx908),
1473  key_type,
1474  value_type,
1475  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1476  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
1477  && (sizeof(value_type) > 2))>>
1478  : radix_sort_onesweep_config<kernel_config<256, 18>,
1479  kernel_config<256, 18>,
1480  8,
1481  block_radix_rank_algorithm::match>
1482 {};
1483 
1484 // Based on key_type = short, value_type = short
1485 template<class key_type, class value_type>
1486 struct default_radix_sort_onesweep_config<
1487  static_cast<unsigned int>(target_arch::gfx908),
1488  key_type,
1489  value_type,
1490  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1491  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
1492  && (sizeof(value_type) > 1))>>
1493  : radix_sort_onesweep_config<kernel_config<1024, 22>,
1494  kernel_config<1024, 22>,
1495  6,
1496  block_radix_rank_algorithm::match>
1497 {};
1498 
1499 // Based on key_type = short, value_type = int8_t
1500 template<class key_type, class value_type>
1501 struct default_radix_sort_onesweep_config<
1502  static_cast<unsigned int>(target_arch::gfx908),
1503  key_type,
1504  value_type,
1505  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1506  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
1507  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1508  : radix_sort_onesweep_config<kernel_config<256, 18>,
1509  kernel_config<256, 18>,
1510  8,
1511  block_radix_rank_algorithm::match>
1512 {};
1513 
1514 // Based on key_type = short, value_type = empty_type
1515 template<class key_type, class value_type>
1516 struct default_radix_sort_onesweep_config<
1517  static_cast<unsigned int>(target_arch::gfx908),
1518  key_type,
1519  value_type,
1520  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1521  && (sizeof(key_type) > 1)
1522  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1523  : radix_sort_onesweep_config<kernel_config<256, 22>,
1524  kernel_config<256, 22>,
1525  8,
1526  block_radix_rank_algorithm::match>
1527 {};
1528 
1529 // Based on key_type = int8_t, value_type = int64_t
1530 template<class key_type, class value_type>
1531 struct default_radix_sort_onesweep_config<
1532  static_cast<unsigned int>(target_arch::gfx908),
1533  key_type,
1534  value_type,
1535  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1536  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1537  : radix_sort_onesweep_config<kernel_config<512, 6>,
1538  kernel_config<512, 6>,
1539  8,
1540  block_radix_rank_algorithm::match>
1541 {};
1542 
1543 // Based on key_type = int8_t, value_type = int
1544 template<class key_type, class value_type>
1545 struct default_radix_sort_onesweep_config<
1546  static_cast<unsigned int>(target_arch::gfx908),
1547  key_type,
1548  value_type,
1549  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1550  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1551  : radix_sort_onesweep_config<kernel_config<256, 22>,
1552  kernel_config<256, 22>,
1553  8,
1554  block_radix_rank_algorithm::match>
1555 {};
1556 
1557 // Based on key_type = int8_t, value_type = short
1558 template<class key_type, class value_type>
1559 struct default_radix_sort_onesweep_config<
1560  static_cast<unsigned int>(target_arch::gfx908),
1561  key_type,
1562  value_type,
1563  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1564  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1565  : radix_sort_onesweep_config<kernel_config<256, 16>,
1566  kernel_config<256, 16>,
1567  8,
1568  block_radix_rank_algorithm::match>
1569 {};
1570 
1571 // Based on key_type = int8_t, value_type = int8_t
1572 template<class key_type, class value_type>
1573 struct default_radix_sort_onesweep_config<
1574  static_cast<unsigned int>(target_arch::gfx908),
1575  key_type,
1576  value_type,
1577  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1578  && (sizeof(value_type) <= 1)
1579  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1580  : radix_sort_onesweep_config<kernel_config<1024, 22>,
1581  kernel_config<1024, 22>,
1582  4,
1583  block_radix_rank_algorithm::basic>
1584 {};
1585 
1586 // Based on key_type = int8_t, value_type = empty_type
1587 template<class key_type, class value_type>
1588 struct default_radix_sort_onesweep_config<
1589  static_cast<unsigned int>(target_arch::gfx908),
1590  key_type,
1591  value_type,
1592  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1593  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1594  : radix_sort_onesweep_config<kernel_config<256, 16>,
1595  kernel_config<256, 16>,
1596  8,
1597  block_radix_rank_algorithm::match>
1598 {};
1599 
1600 // Based on key_type = double, value_type = int64_t
1601 template<class key_type, class value_type>
1602 struct default_radix_sort_onesweep_config<
1603  static_cast<unsigned int>(target_arch::unknown),
1604  key_type,
1605  value_type,
1606  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1607  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1608  && (sizeof(value_type) > 4))>>
1609  : radix_sort_onesweep_config<kernel_config<1024, 6>,
1610  kernel_config<1024, 6>,
1611  8,
1612  block_radix_rank_algorithm::match>
1613 {};
1614 
1615 // Based on key_type = double, value_type = int
1616 template<class key_type, class value_type>
1617 struct default_radix_sort_onesweep_config<
1618  static_cast<unsigned int>(target_arch::unknown),
1619  key_type,
1620  value_type,
1621  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1622  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1623  && (sizeof(value_type) > 2))>>
1624  : radix_sort_onesweep_config<kernel_config<256, 12>,
1625  kernel_config<256, 12>,
1626  8,
1627  block_radix_rank_algorithm::match>
1628 {};
1629 
1630 // Based on key_type = double, value_type = short
1631 template<class key_type, class value_type>
1632 struct default_radix_sort_onesweep_config<
1633  static_cast<unsigned int>(target_arch::unknown),
1634  key_type,
1635  value_type,
1636  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1637  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
1638  && (sizeof(value_type) > 1))>>
1639  : radix_sort_onesweep_config<kernel_config<256, 12>,
1640  kernel_config<256, 12>,
1641  8,
1642  block_radix_rank_algorithm::match>
1643 {};
1644 
1645 // Based on key_type = double, value_type = int8_t
1646 template<class key_type, class value_type>
1647 struct default_radix_sort_onesweep_config<
1648  static_cast<unsigned int>(target_arch::unknown),
1649  key_type,
1650  value_type,
1651  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1652  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
1653  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1654  : radix_sort_onesweep_config<kernel_config<512, 6>,
1655  kernel_config<512, 6>,
1656  8,
1657  block_radix_rank_algorithm::match>
1658 {};
1659 
1660 // Based on key_type = double, value_type = empty_type
1661 template<class key_type, class value_type>
1662 struct default_radix_sort_onesweep_config<
1663  static_cast<unsigned int>(target_arch::unknown),
1664  key_type,
1665  value_type,
1666  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1667  && (sizeof(key_type) > 4)
1668  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1669  : radix_sort_onesweep_config<kernel_config<512, 6>,
1670  kernel_config<512, 6>,
1671  8,
1672  block_radix_rank_algorithm::match>
1673 {};
1674 
1675 // Based on key_type = float, value_type = int64_t
1676 template<class key_type, class value_type>
1677 struct default_radix_sort_onesweep_config<
1678  static_cast<unsigned int>(target_arch::unknown),
1679  key_type,
1680  value_type,
1681  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1682  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
1683  && (sizeof(value_type) > 4))>>
1684  : radix_sort_onesweep_config<kernel_config<1024, 6>,
1685  kernel_config<1024, 6>,
1686  8,
1687  block_radix_rank_algorithm::match>
1688 {};
1689 
1690 // Based on key_type = float, value_type = int
1691 template<class key_type, class value_type>
1692 struct default_radix_sort_onesweep_config<
1693  static_cast<unsigned int>(target_arch::unknown),
1694  key_type,
1695  value_type,
1696  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1697  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1698  && (sizeof(value_type) > 2))>>
1699  : radix_sort_onesweep_config<kernel_config<256, 22>,
1700  kernel_config<256, 22>,
1701  8,
1702  block_radix_rank_algorithm::match>
1703 {};
1704 
1705 // Based on key_type = float, value_type = short
1706 template<class key_type, class value_type>
1707 struct default_radix_sort_onesweep_config<
1708  static_cast<unsigned int>(target_arch::unknown),
1709  key_type,
1710  value_type,
1711  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1712  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1713  && (sizeof(value_type) > 1))>>
1714  : radix_sort_onesweep_config<kernel_config<256, 22>,
1715  kernel_config<256, 22>,
1716  8,
1717  block_radix_rank_algorithm::match>
1718 {};
1719 
1720 // Based on key_type = float, value_type = int8_t
1721 template<class key_type, class value_type>
1722 struct default_radix_sort_onesweep_config<
1723  static_cast<unsigned int>(target_arch::unknown),
1724  key_type,
1725  value_type,
1726  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1727  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
1728  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1729  : radix_sort_onesweep_config<kernel_config<256, 22>,
1730  kernel_config<256, 22>,
1731  8,
1732  block_radix_rank_algorithm::match>
1733 {};
1734 
1735 // Based on key_type = float, value_type = empty_type
1736 template<class key_type, class value_type>
1737 struct default_radix_sort_onesweep_config<
1738  static_cast<unsigned int>(target_arch::unknown),
1739  key_type,
1740  value_type,
1741  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1742  && (sizeof(key_type) > 2)
1743  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1744  : radix_sort_onesweep_config<kernel_config<256, 22>,
1745  kernel_config<256, 22>,
1746  8,
1747  block_radix_rank_algorithm::match>
1748 {};
1749 
1750 // Based on key_type = rocprim::half, value_type = int64_t
1751 template<class key_type, class value_type>
1752 struct default_radix_sort_onesweep_config<
1753  static_cast<unsigned int>(target_arch::unknown),
1754  key_type,
1755  value_type,
1756  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1757  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1758  : radix_sort_onesweep_config<kernel_config<256, 12>,
1759  kernel_config<256, 12>,
1760  8,
1761  block_radix_rank_algorithm::match>
1762 {};
1763 
1764 // Based on key_type = rocprim::half, value_type = int
1765 template<class key_type, class value_type>
1766 struct default_radix_sort_onesweep_config<
1767  static_cast<unsigned int>(target_arch::unknown),
1768  key_type,
1769  value_type,
1770  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1771  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1772  : radix_sort_onesweep_config<kernel_config<256, 16>,
1773  kernel_config<256, 16>,
1774  8,
1775  block_radix_rank_algorithm::match>
1776 {};
1777 
1778 // Based on key_type = rocprim::half, value_type = short
1779 template<class key_type, class value_type>
1780 struct default_radix_sort_onesweep_config<
1781  static_cast<unsigned int>(target_arch::unknown),
1782  key_type,
1783  value_type,
1784  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1785  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1786  : radix_sort_onesweep_config<kernel_config<256, 16>,
1787  kernel_config<256, 16>,
1788  8,
1789  block_radix_rank_algorithm::match>
1790 {};
1791 
1792 // Based on key_type = rocprim::half, value_type = int8_t
1793 template<class key_type, class value_type>
1794 struct default_radix_sort_onesweep_config<
1795  static_cast<unsigned int>(target_arch::unknown),
1796  key_type,
1797  value_type,
1798  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1799  && (sizeof(value_type) <= 1)
1800  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1801  : radix_sort_onesweep_config<kernel_config<256, 16>,
1802  kernel_config<256, 16>,
1803  8,
1804  block_radix_rank_algorithm::match>
1805 {};
1806 
1807 // Based on key_type = rocprim::half, value_type = empty_type
1808 template<class key_type, class value_type>
1809 struct default_radix_sort_onesweep_config<
1810  static_cast<unsigned int>(target_arch::unknown),
1811  key_type,
1812  value_type,
1813  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1814  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1815  : radix_sort_onesweep_config<kernel_config<256, 16>,
1816  kernel_config<256, 16>,
1817  8,
1818  block_radix_rank_algorithm::match>
1819 {};
1820 
1821 // Based on key_type = int64_t, value_type = int64_t
1822 template<class key_type, class value_type>
1823 struct default_radix_sort_onesweep_config<
1824  static_cast<unsigned int>(target_arch::unknown),
1825  key_type,
1826  value_type,
1827  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1828  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1829  && (sizeof(value_type) > 4))>>
1830  : radix_sort_onesweep_config<kernel_config<1024, 6>,
1831  kernel_config<1024, 6>,
1832  6,
1833  block_radix_rank_algorithm::match>
1834 {};
1835 
1836 // Based on key_type = int64_t, value_type = int
1837 template<class key_type, class value_type>
1838 struct default_radix_sort_onesweep_config<
1839  static_cast<unsigned int>(target_arch::unknown),
1840  key_type,
1841  value_type,
1842  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1843  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1844  && (sizeof(value_type) > 2))>>
1845  : radix_sort_onesweep_config<kernel_config<256, 12>,
1846  kernel_config<256, 12>,
1847  8,
1848  block_radix_rank_algorithm::match>
1849 {};
1850 
1851 // Based on key_type = int64_t, value_type = short
1852 template<class key_type, class value_type>
1853 struct default_radix_sort_onesweep_config<
1854  static_cast<unsigned int>(target_arch::unknown),
1855  key_type,
1856  value_type,
1857  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1858  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
1859  && (sizeof(value_type) > 1))>>
1860  : radix_sort_onesweep_config<kernel_config<256, 12>,
1861  kernel_config<256, 12>,
1862  8,
1863  block_radix_rank_algorithm::match>
1864 {};
1865 
1866 // Based on key_type = int64_t, value_type = int8_t
1867 template<class key_type, class value_type>
1868 struct default_radix_sort_onesweep_config<
1869  static_cast<unsigned int>(target_arch::unknown),
1870  key_type,
1871  value_type,
1872  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1873  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
1874  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1875  : radix_sort_onesweep_config<kernel_config<512, 6>,
1876  kernel_config<512, 6>,
1877  8,
1878  block_radix_rank_algorithm::match>
1879 {};
1880 
1881 // Based on key_type = int64_t, value_type = empty_type
1882 template<class key_type, class value_type>
1883 struct default_radix_sort_onesweep_config<
1884  static_cast<unsigned int>(target_arch::unknown),
1885  key_type,
1886  value_type,
1887  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1888  && (sizeof(key_type) > 4)
1889  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1890  : radix_sort_onesweep_config<kernel_config<256, 12>,
1891  kernel_config<256, 12>,
1892  8,
1893  block_radix_rank_algorithm::match>
1894 {};
1895 
1896 // Based on key_type = int, value_type = int64_t
1897 template<class key_type, class value_type>
1898 struct default_radix_sort_onesweep_config<
1899  static_cast<unsigned int>(target_arch::unknown),
1900  key_type,
1901  value_type,
1902  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1903  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
1904  && (sizeof(value_type) > 4))>>
1905  : radix_sort_onesweep_config<kernel_config<256, 12>,
1906  kernel_config<256, 12>,
1907  8,
1908  block_radix_rank_algorithm::match>
1909 {};
1910 
1911 // Based on key_type = int, value_type = int
1912 template<class key_type, class value_type>
1913 struct default_radix_sort_onesweep_config<
1914  static_cast<unsigned int>(target_arch::unknown),
1915  key_type,
1916  value_type,
1917  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1918  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1919  && (sizeof(value_type) > 2))>>
1920  : radix_sort_onesweep_config<kernel_config<512, 12>,
1921  kernel_config<512, 12>,
1922  5,
1923  block_radix_rank_algorithm::match>
1924 {};
1925 
1926 // Based on key_type = int, value_type = short
1927 template<class key_type, class value_type>
1928 struct default_radix_sort_onesweep_config<
1929  static_cast<unsigned int>(target_arch::unknown),
1930  key_type,
1931  value_type,
1932  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1933  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1934  && (sizeof(value_type) > 1))>>
1935  : radix_sort_onesweep_config<kernel_config<256, 16>,
1936  kernel_config<256, 16>,
1937  8,
1938  block_radix_rank_algorithm::match>
1939 {};
1940 
1941 // Based on key_type = int, value_type = int8_t
1942 template<class key_type, class value_type>
1943 struct default_radix_sort_onesweep_config<
1944  static_cast<unsigned int>(target_arch::unknown),
1945  key_type,
1946  value_type,
1947  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1948  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
1949  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
1950  : radix_sort_onesweep_config<kernel_config<256, 22>,
1951  kernel_config<256, 22>,
1952  8,
1953  block_radix_rank_algorithm::match>
1954 {};
1955 
1956 // Based on key_type = int, value_type = empty_type
1957 template<class key_type, class value_type>
1958 struct default_radix_sort_onesweep_config<
1959  static_cast<unsigned int>(target_arch::unknown),
1960  key_type,
1961  value_type,
1962  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1963  && (sizeof(key_type) > 2)
1964  && (std::is_same<value_type, rocprim::empty_type>::value))>>
1965  : radix_sort_onesweep_config<kernel_config<256, 22>,
1966  kernel_config<256, 22>,
1967  7,
1968  block_radix_rank_algorithm::match>
1969 {};
1970 
1971 // Based on key_type = short, value_type = int64_t
1972 template<class key_type, class value_type>
1973 struct default_radix_sort_onesweep_config<
1974  static_cast<unsigned int>(target_arch::unknown),
1975  key_type,
1976  value_type,
1977  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1978  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
1979  && (sizeof(value_type) > 4))>>
1980  : radix_sort_onesweep_config<kernel_config<256, 18>,
1981  kernel_config<256, 18>,
1982  8,
1983  block_radix_rank_algorithm::match>
1984 {};
1985 
1986 // Based on key_type = short, value_type = int
1987 template<class key_type, class value_type>
1988 struct default_radix_sort_onesweep_config<
1989  static_cast<unsigned int>(target_arch::unknown),
1990  key_type,
1991  value_type,
1992  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1993  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
1994  && (sizeof(value_type) > 2))>>
1995  : radix_sort_onesweep_config<kernel_config<256, 18>,
1996  kernel_config<256, 18>,
1997  8,
1998  block_radix_rank_algorithm::match>
1999 {};
2000 
2001 // Based on key_type = short, value_type = short
2002 template<class key_type, class value_type>
2003 struct default_radix_sort_onesweep_config<
2004  static_cast<unsigned int>(target_arch::unknown),
2005  key_type,
2006  value_type,
2007  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2008  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
2009  && (sizeof(value_type) > 1))>>
2010  : radix_sort_onesweep_config<kernel_config<1024, 22>,
2011  kernel_config<1024, 22>,
2012  6,
2013  block_radix_rank_algorithm::match>
2014 {};
2015 
2016 // Based on key_type = short, value_type = int8_t
2017 template<class key_type, class value_type>
2018 struct default_radix_sort_onesweep_config<
2019  static_cast<unsigned int>(target_arch::unknown),
2020  key_type,
2021  value_type,
2022  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2023  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
2024  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2025  : radix_sort_onesweep_config<kernel_config<256, 18>,
2026  kernel_config<256, 18>,
2027  8,
2028  block_radix_rank_algorithm::match>
2029 {};
2030 
2031 // Based on key_type = short, value_type = empty_type
2032 template<class key_type, class value_type>
2033 struct default_radix_sort_onesweep_config<
2034  static_cast<unsigned int>(target_arch::unknown),
2035  key_type,
2036  value_type,
2037  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2038  && (sizeof(key_type) > 1)
2039  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2040  : radix_sort_onesweep_config<kernel_config<256, 22>,
2041  kernel_config<256, 22>,
2042  8,
2043  block_radix_rank_algorithm::match>
2044 {};
2045 
2046 // Based on key_type = int8_t, value_type = int64_t
2047 template<class key_type, class value_type>
2048 struct default_radix_sort_onesweep_config<
2049  static_cast<unsigned int>(target_arch::unknown),
2050  key_type,
2051  value_type,
2052  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2053  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
2054  : radix_sort_onesweep_config<kernel_config<512, 6>,
2055  kernel_config<512, 6>,
2056  8,
2057  block_radix_rank_algorithm::match>
2058 {};
2059 
2060 // Based on key_type = int8_t, value_type = int
2061 template<class key_type, class value_type>
2062 struct default_radix_sort_onesweep_config<
2063  static_cast<unsigned int>(target_arch::unknown),
2064  key_type,
2065  value_type,
2066  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2067  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
2068  : radix_sort_onesweep_config<kernel_config<256, 22>,
2069  kernel_config<256, 22>,
2070  8,
2071  block_radix_rank_algorithm::match>
2072 {};
2073 
2074 // Based on key_type = int8_t, value_type = short
2075 template<class key_type, class value_type>
2076 struct default_radix_sort_onesweep_config<
2077  static_cast<unsigned int>(target_arch::unknown),
2078  key_type,
2079  value_type,
2080  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2081  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
2082  : radix_sort_onesweep_config<kernel_config<256, 16>,
2083  kernel_config<256, 16>,
2084  8,
2085  block_radix_rank_algorithm::match>
2086 {};
2087 
2088 // Based on key_type = int8_t, value_type = int8_t
2089 template<class key_type, class value_type>
2090 struct default_radix_sort_onesweep_config<
2091  static_cast<unsigned int>(target_arch::unknown),
2092  key_type,
2093  value_type,
2094  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2095  && (sizeof(value_type) <= 1)
2096  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2097  : radix_sort_onesweep_config<kernel_config<1024, 22>,
2098  kernel_config<1024, 22>,
2099  4,
2100  block_radix_rank_algorithm::basic>
2101 {};
2102 
2103 // Based on key_type = int8_t, value_type = empty_type
2104 template<class key_type, class value_type>
2105 struct default_radix_sort_onesweep_config<
2106  static_cast<unsigned int>(target_arch::unknown),
2107  key_type,
2108  value_type,
2109  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2110  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2111  : radix_sort_onesweep_config<kernel_config<256, 16>,
2112  kernel_config<256, 16>,
2113  8,
2114  block_radix_rank_algorithm::match>
2115 {};
2116 
2117 // Based on key_type = double, value_type = int64_t
2118 template<class key_type, class value_type>
2119 struct default_radix_sort_onesweep_config<
2120  static_cast<unsigned int>(target_arch::gfx90a),
2121  key_type,
2122  value_type,
2123  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2124  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
2125  && (sizeof(value_type) > 4))>>
2126  : radix_sort_onesweep_config<kernel_config<1024, 6>,
2127  kernel_config<1024, 6>,
2128  8,
2129  block_radix_rank_algorithm::match>
2130 {};
2131 
2132 // Based on key_type = double, value_type = int
2133 template<class key_type, class value_type>
2134 struct default_radix_sort_onesweep_config<
2135  static_cast<unsigned int>(target_arch::gfx90a),
2136  key_type,
2137  value_type,
2138  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2139  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
2140  && (sizeof(value_type) > 2))>>
2141  : radix_sort_onesweep_config<kernel_config<256, 12>,
2142  kernel_config<256, 12>,
2143  8,
2144  block_radix_rank_algorithm::match>
2145 {};
2146 
2147 // Based on key_type = double, value_type = short
2148 template<class key_type, class value_type>
2149 struct default_radix_sort_onesweep_config<
2150  static_cast<unsigned int>(target_arch::gfx90a),
2151  key_type,
2152  value_type,
2153  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2154  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
2155  && (sizeof(value_type) > 1))>>
2156  : radix_sort_onesweep_config<kernel_config<256, 12>,
2157  kernel_config<256, 12>,
2158  8,
2159  block_radix_rank_algorithm::match>
2160 {};
2161 
2162 // Based on key_type = double, value_type = int8_t
2163 template<class key_type, class value_type>
2164 struct default_radix_sort_onesweep_config<
2165  static_cast<unsigned int>(target_arch::gfx90a),
2166  key_type,
2167  value_type,
2168  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2169  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
2170  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2171  : radix_sort_onesweep_config<kernel_config<512, 6>,
2172  kernel_config<512, 6>,
2173  8,
2174  block_radix_rank_algorithm::match>
2175 {};
2176 
2177 // Based on key_type = double, value_type = empty_type
2178 template<class key_type, class value_type>
2179 struct default_radix_sort_onesweep_config<
2180  static_cast<unsigned int>(target_arch::gfx90a),
2181  key_type,
2182  value_type,
2183  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2184  && (sizeof(key_type) > 4)
2185  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2186  : radix_sort_onesweep_config<kernel_config<512, 6>,
2187  kernel_config<512, 6>,
2188  8,
2189  block_radix_rank_algorithm::match>
2190 {};
2191 
2192 // Based on key_type = float, value_type = int64_t
2193 template<class key_type, class value_type>
2194 struct default_radix_sort_onesweep_config<
2195  static_cast<unsigned int>(target_arch::gfx90a),
2196  key_type,
2197  value_type,
2198  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2199  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
2200  && (sizeof(value_type) > 4))>>
2201  : radix_sort_onesweep_config<kernel_config<1024, 6>,
2202  kernel_config<1024, 6>,
2203  8,
2204  block_radix_rank_algorithm::match>
2205 {};
2206 
2207 // Based on key_type = float, value_type = int
2208 template<class key_type, class value_type>
2209 struct default_radix_sort_onesweep_config<
2210  static_cast<unsigned int>(target_arch::gfx90a),
2211  key_type,
2212  value_type,
2213  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2214  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
2215  && (sizeof(value_type) > 2))>>
2216  : radix_sort_onesweep_config<kernel_config<256, 22>,
2217  kernel_config<256, 22>,
2218  8,
2219  block_radix_rank_algorithm::match>
2220 {};
2221 
2222 // Based on key_type = float, value_type = short
2223 template<class key_type, class value_type>
2224 struct default_radix_sort_onesweep_config<
2225  static_cast<unsigned int>(target_arch::gfx90a),
2226  key_type,
2227  value_type,
2228  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2229  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
2230  && (sizeof(value_type) > 1))>>
2231  : radix_sort_onesweep_config<kernel_config<256, 22>,
2232  kernel_config<256, 22>,
2233  8,
2234  block_radix_rank_algorithm::match>
2235 {};
2236 
2237 // Based on key_type = float, value_type = int8_t
2238 template<class key_type, class value_type>
2239 struct default_radix_sort_onesweep_config<
2240  static_cast<unsigned int>(target_arch::gfx90a),
2241  key_type,
2242  value_type,
2243  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2244  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
2245  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2246  : radix_sort_onesweep_config<kernel_config<256, 22>,
2247  kernel_config<256, 22>,
2248  8,
2249  block_radix_rank_algorithm::match>
2250 {};
2251 
2252 // Based on key_type = float, value_type = empty_type
2253 template<class key_type, class value_type>
2254 struct default_radix_sort_onesweep_config<
2255  static_cast<unsigned int>(target_arch::gfx90a),
2256  key_type,
2257  value_type,
2258  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2259  && (sizeof(key_type) > 2)
2260  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2261  : radix_sort_onesweep_config<kernel_config<256, 22>,
2262  kernel_config<256, 22>,
2263  8,
2264  block_radix_rank_algorithm::match>
2265 {};
2266 
2267 // Based on key_type = rocprim::half, value_type = int64_t
2268 template<class key_type, class value_type>
2269 struct default_radix_sort_onesweep_config<
2270  static_cast<unsigned int>(target_arch::gfx90a),
2271  key_type,
2272  value_type,
2273  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2274  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
2275  : radix_sort_onesweep_config<kernel_config<256, 12>,
2276  kernel_config<256, 12>,
2277  8,
2278  block_radix_rank_algorithm::match>
2279 {};
2280 
2281 // Based on key_type = rocprim::half, value_type = int
2282 template<class key_type, class value_type>
2283 struct default_radix_sort_onesweep_config<
2284  static_cast<unsigned int>(target_arch::gfx90a),
2285  key_type,
2286  value_type,
2287  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2288  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
2289  : radix_sort_onesweep_config<kernel_config<256, 16>,
2290  kernel_config<256, 16>,
2291  8,
2292  block_radix_rank_algorithm::match>
2293 {};
2294 
2295 // Based on key_type = rocprim::half, value_type = short
2296 template<class key_type, class value_type>
2297 struct default_radix_sort_onesweep_config<
2298  static_cast<unsigned int>(target_arch::gfx90a),
2299  key_type,
2300  value_type,
2301  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2302  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
2303  : radix_sort_onesweep_config<kernel_config<256, 16>,
2304  kernel_config<256, 16>,
2305  8,
2306  block_radix_rank_algorithm::match>
2307 {};
2308 
2309 // Based on key_type = rocprim::half, value_type = int8_t
2310 template<class key_type, class value_type>
2311 struct default_radix_sort_onesweep_config<
2312  static_cast<unsigned int>(target_arch::gfx90a),
2313  key_type,
2314  value_type,
2315  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2316  && (sizeof(value_type) <= 1)
2317  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2318  : radix_sort_onesweep_config<kernel_config<256, 16>,
2319  kernel_config<256, 16>,
2320  8,
2321  block_radix_rank_algorithm::match>
2322 {};
2323 
2324 // Based on key_type = rocprim::half, value_type = empty_type
2325 template<class key_type, class value_type>
2326 struct default_radix_sort_onesweep_config<
2327  static_cast<unsigned int>(target_arch::gfx90a),
2328  key_type,
2329  value_type,
2330  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2331  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2332  : radix_sort_onesweep_config<kernel_config<256, 16>,
2333  kernel_config<256, 16>,
2334  8,
2335  block_radix_rank_algorithm::match>
2336 {};
2337 
2338 // Based on key_type = int64_t, value_type = int64_t
2339 template<class key_type, class value_type>
2340 struct default_radix_sort_onesweep_config<
2341  static_cast<unsigned int>(target_arch::gfx90a),
2342  key_type,
2343  value_type,
2344  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2345  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
2346  && (sizeof(value_type) > 4))>>
2347  : radix_sort_onesweep_config<kernel_config<1024, 6>,
2348  kernel_config<1024, 6>,
2349  6,
2350  block_radix_rank_algorithm::match>
2351 {};
2352 
2353 // Based on key_type = int64_t, value_type = int
2354 template<class key_type, class value_type>
2355 struct default_radix_sort_onesweep_config<
2356  static_cast<unsigned int>(target_arch::gfx90a),
2357  key_type,
2358  value_type,
2359  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2360  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
2361  && (sizeof(value_type) > 2))>>
2362  : radix_sort_onesweep_config<kernel_config<256, 12>,
2363  kernel_config<256, 12>,
2364  8,
2365  block_radix_rank_algorithm::match>
2366 {};
2367 
2368 // Based on key_type = int64_t, value_type = short
2369 template<class key_type, class value_type>
2370 struct default_radix_sort_onesweep_config<
2371  static_cast<unsigned int>(target_arch::gfx90a),
2372  key_type,
2373  value_type,
2374  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2375  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
2376  && (sizeof(value_type) > 1))>>
2377  : radix_sort_onesweep_config<kernel_config<256, 12>,
2378  kernel_config<256, 12>,
2379  8,
2380  block_radix_rank_algorithm::match>
2381 {};
2382 
2383 // Based on key_type = int64_t, value_type = int8_t
2384 template<class key_type, class value_type>
2385 struct default_radix_sort_onesweep_config<
2386  static_cast<unsigned int>(target_arch::gfx90a),
2387  key_type,
2388  value_type,
2389  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2390  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
2391  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2392  : radix_sort_onesweep_config<kernel_config<512, 6>,
2393  kernel_config<512, 6>,
2394  8,
2395  block_radix_rank_algorithm::match>
2396 {};
2397 
2398 // Based on key_type = int64_t, value_type = empty_type
2399 template<class key_type, class value_type>
2400 struct default_radix_sort_onesweep_config<
2401  static_cast<unsigned int>(target_arch::gfx90a),
2402  key_type,
2403  value_type,
2404  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2405  && (sizeof(key_type) > 4)
2406  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2407  : radix_sort_onesweep_config<kernel_config<256, 12>,
2408  kernel_config<256, 12>,
2409  8,
2410  block_radix_rank_algorithm::match>
2411 {};
2412 
2413 // Based on key_type = int, value_type = int64_t
2414 template<class key_type, class value_type>
2415 struct default_radix_sort_onesweep_config<
2416  static_cast<unsigned int>(target_arch::gfx90a),
2417  key_type,
2418  value_type,
2419  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2420  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
2421  && (sizeof(value_type) > 4))>>
2422  : radix_sort_onesweep_config<kernel_config<256, 12>,
2423  kernel_config<256, 12>,
2424  8,
2425  block_radix_rank_algorithm::match>
2426 {};
2427 
2428 // Based on key_type = int, value_type = int
2429 template<class key_type, class value_type>
2430 struct default_radix_sort_onesweep_config<
2431  static_cast<unsigned int>(target_arch::gfx90a),
2432  key_type,
2433  value_type,
2434  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2435  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
2436  && (sizeof(value_type) > 2))>>
2437  : radix_sort_onesweep_config<kernel_config<512, 12>,
2438  kernel_config<512, 12>,
2439  5,
2440  block_radix_rank_algorithm::match>
2441 {};
2442 
2443 // Based on key_type = int, value_type = short
2444 template<class key_type, class value_type>
2445 struct default_radix_sort_onesweep_config<
2446  static_cast<unsigned int>(target_arch::gfx90a),
2447  key_type,
2448  value_type,
2449  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2450  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
2451  && (sizeof(value_type) > 1))>>
2452  : radix_sort_onesweep_config<kernel_config<256, 16>,
2453  kernel_config<256, 16>,
2454  8,
2455  block_radix_rank_algorithm::match>
2456 {};
2457 
2458 // Based on key_type = int, value_type = int8_t
2459 template<class key_type, class value_type>
2460 struct default_radix_sort_onesweep_config<
2461  static_cast<unsigned int>(target_arch::gfx90a),
2462  key_type,
2463  value_type,
2464  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2465  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
2466  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2467  : radix_sort_onesweep_config<kernel_config<256, 22>,
2468  kernel_config<256, 22>,
2469  8,
2470  block_radix_rank_algorithm::match>
2471 {};
2472 
2473 // Based on key_type = int, value_type = empty_type
2474 template<class key_type, class value_type>
2475 struct default_radix_sort_onesweep_config<
2476  static_cast<unsigned int>(target_arch::gfx90a),
2477  key_type,
2478  value_type,
2479  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2480  && (sizeof(key_type) > 2)
2481  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2482  : radix_sort_onesweep_config<kernel_config<256, 22>,
2483  kernel_config<256, 22>,
2484  7,
2485  block_radix_rank_algorithm::match>
2486 {};
2487 
2488 // Based on key_type = short, value_type = int64_t
2489 template<class key_type, class value_type>
2490 struct default_radix_sort_onesweep_config<
2491  static_cast<unsigned int>(target_arch::gfx90a),
2492  key_type,
2493  value_type,
2494  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2495  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
2496  && (sizeof(value_type) > 4))>>
2497  : radix_sort_onesweep_config<kernel_config<256, 18>,
2498  kernel_config<256, 18>,
2499  8,
2500  block_radix_rank_algorithm::match>
2501 {};
2502 
2503 // Based on key_type = short, value_type = int
2504 template<class key_type, class value_type>
2505 struct default_radix_sort_onesweep_config<
2506  static_cast<unsigned int>(target_arch::gfx90a),
2507  key_type,
2508  value_type,
2509  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2510  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
2511  && (sizeof(value_type) > 2))>>
2512  : radix_sort_onesweep_config<kernel_config<256, 18>,
2513  kernel_config<256, 18>,
2514  8,
2515  block_radix_rank_algorithm::match>
2516 {};
2517 
2518 // Based on key_type = short, value_type = short
2519 template<class key_type, class value_type>
2520 struct default_radix_sort_onesweep_config<
2521  static_cast<unsigned int>(target_arch::gfx90a),
2522  key_type,
2523  value_type,
2524  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2525  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
2526  && (sizeof(value_type) > 1))>>
2527  : radix_sort_onesweep_config<kernel_config<1024, 22>,
2528  kernel_config<1024, 22>,
2529  6,
2530  block_radix_rank_algorithm::match>
2531 {};
2532 
2533 // Based on key_type = short, value_type = int8_t
2534 template<class key_type, class value_type>
2535 struct default_radix_sort_onesweep_config<
2536  static_cast<unsigned int>(target_arch::gfx90a),
2537  key_type,
2538  value_type,
2539  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2540  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
2541  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2542  : radix_sort_onesweep_config<kernel_config<256, 18>,
2543  kernel_config<256, 18>,
2544  8,
2545  block_radix_rank_algorithm::match>
2546 {};
2547 
2548 // Based on key_type = short, value_type = empty_type
2549 template<class key_type, class value_type>
2550 struct default_radix_sort_onesweep_config<
2551  static_cast<unsigned int>(target_arch::gfx90a),
2552  key_type,
2553  value_type,
2554  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2555  && (sizeof(key_type) > 1)
2556  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2557  : radix_sort_onesweep_config<kernel_config<256, 22>,
2558  kernel_config<256, 22>,
2559  8,
2560  block_radix_rank_algorithm::match>
2561 {};
2562 
2563 // Based on key_type = int8_t, value_type = int64_t
2564 template<class key_type, class value_type>
2565 struct default_radix_sort_onesweep_config<
2566  static_cast<unsigned int>(target_arch::gfx90a),
2567  key_type,
2568  value_type,
2569  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2570  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
2571  : radix_sort_onesweep_config<kernel_config<512, 6>,
2572  kernel_config<512, 6>,
2573  8,
2574  block_radix_rank_algorithm::match>
2575 {};
2576 
2577 // Based on key_type = int8_t, value_type = int
2578 template<class key_type, class value_type>
2579 struct default_radix_sort_onesweep_config<
2580  static_cast<unsigned int>(target_arch::gfx90a),
2581  key_type,
2582  value_type,
2583  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2584  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
2585  : radix_sort_onesweep_config<kernel_config<256, 22>,
2586  kernel_config<256, 22>,
2587  8,
2588  block_radix_rank_algorithm::match>
2589 {};
2590 
2591 // Based on key_type = int8_t, value_type = short
2592 template<class key_type, class value_type>
2593 struct default_radix_sort_onesweep_config<
2594  static_cast<unsigned int>(target_arch::gfx90a),
2595  key_type,
2596  value_type,
2597  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2598  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
2599  : radix_sort_onesweep_config<kernel_config<256, 16>,
2600  kernel_config<256, 16>,
2601  8,
2602  block_radix_rank_algorithm::match>
2603 {};
2604 
2605 // Based on key_type = int8_t, value_type = int8_t
2606 template<class key_type, class value_type>
2607 struct default_radix_sort_onesweep_config<
2608  static_cast<unsigned int>(target_arch::gfx90a),
2609  key_type,
2610  value_type,
2611  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2612  && (sizeof(value_type) <= 1)
2613  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2614  : radix_sort_onesweep_config<kernel_config<1024, 22>,
2615  kernel_config<1024, 22>,
2616  4,
2617  block_radix_rank_algorithm::basic>
2618 {};
2619 
2620 // Based on key_type = int8_t, value_type = empty_type
2621 template<class key_type, class value_type>
2622 struct default_radix_sort_onesweep_config<
2623  static_cast<unsigned int>(target_arch::gfx90a),
2624  key_type,
2625  value_type,
2626  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2627  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2628  : radix_sort_onesweep_config<kernel_config<256, 16>,
2629  kernel_config<256, 16>,
2630  8,
2631  block_radix_rank_algorithm::match>
2632 {};
2633 
2634 // Based on key_type = double, value_type = int64_t
2635 template<class key_type, class value_type>
2636 struct default_radix_sort_onesweep_config<
2637  static_cast<unsigned int>(target_arch::gfx906),
2638  key_type,
2639  value_type,
2640  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2641  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
2642  && (sizeof(value_type) > 4))>>
2643  : radix_sort_onesweep_config<kernel_config<256, 12>,
2644  kernel_config<256, 12>,
2645  8,
2646  block_radix_rank_algorithm::match>
2647 {};
2648 
2649 // Based on key_type = double, value_type = int
2650 template<class key_type, class value_type>
2651 struct default_radix_sort_onesweep_config<
2652  static_cast<unsigned int>(target_arch::gfx906),
2653  key_type,
2654  value_type,
2655  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2656  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
2657  && (sizeof(value_type) > 2))>>
2658  : radix_sort_onesweep_config<kernel_config<256, 12>,
2659  kernel_config<256, 12>,
2660  8,
2661  block_radix_rank_algorithm::match>
2662 {};
2663 
2664 // Based on key_type = double, value_type = short
2665 template<class key_type, class value_type>
2666 struct default_radix_sort_onesweep_config<
2667  static_cast<unsigned int>(target_arch::gfx906),
2668  key_type,
2669  value_type,
2670  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2671  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
2672  && (sizeof(value_type) > 1))>>
2673  : radix_sort_onesweep_config<kernel_config<512, 6>,
2674  kernel_config<512, 6>,
2675  8,
2676  block_radix_rank_algorithm::match>
2677 {};
2678 
2679 // Based on key_type = double, value_type = int8_t
2680 template<class key_type, class value_type>
2681 struct default_radix_sort_onesweep_config<
2682  static_cast<unsigned int>(target_arch::gfx906),
2683  key_type,
2684  value_type,
2685  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2686  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
2687  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2688  : radix_sort_onesweep_config<kernel_config<512, 6>,
2689  kernel_config<512, 6>,
2690  8,
2691  block_radix_rank_algorithm::match>
2692 {};
2693 
2694 // Based on key_type = double, value_type = empty_type
2695 template<class key_type, class value_type>
2696 struct default_radix_sort_onesweep_config<
2697  static_cast<unsigned int>(target_arch::gfx906),
2698  key_type,
2699  value_type,
2700  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2701  && (sizeof(key_type) > 4)
2702  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2703  : radix_sort_onesweep_config<kernel_config<256, 12>,
2704  kernel_config<256, 12>,
2705  8,
2706  block_radix_rank_algorithm::match>
2707 {};
2708 
2709 // Based on key_type = float, value_type = int64_t
2710 template<class key_type, class value_type>
2711 struct default_radix_sort_onesweep_config<
2712  static_cast<unsigned int>(target_arch::gfx906),
2713  key_type,
2714  value_type,
2715  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2716  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
2717  && (sizeof(value_type) > 4))>>
2718  : radix_sort_onesweep_config<kernel_config<256, 12>,
2719  kernel_config<256, 12>,
2720  8,
2721  block_radix_rank_algorithm::match>
2722 {};
2723 
2724 // Based on key_type = float, value_type = int
2725 template<class key_type, class value_type>
2726 struct default_radix_sort_onesweep_config<
2727  static_cast<unsigned int>(target_arch::gfx906),
2728  key_type,
2729  value_type,
2730  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2731  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
2732  && (sizeof(value_type) > 2))>>
2733  : radix_sort_onesweep_config<kernel_config<256, 22>,
2734  kernel_config<256, 22>,
2735  8,
2736  block_radix_rank_algorithm::match>
2737 {};
2738 
2739 // Based on key_type = float, value_type = short
2740 template<class key_type, class value_type>
2741 struct default_radix_sort_onesweep_config<
2742  static_cast<unsigned int>(target_arch::gfx906),
2743  key_type,
2744  value_type,
2745  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2746  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
2747  && (sizeof(value_type) > 1))>>
2748  : radix_sort_onesweep_config<kernel_config<256, 16>,
2749  kernel_config<256, 16>,
2750  8,
2751  block_radix_rank_algorithm::match>
2752 {};
2753 
2754 // Based on key_type = float, value_type = int8_t
2755 template<class key_type, class value_type>
2756 struct default_radix_sort_onesweep_config<
2757  static_cast<unsigned int>(target_arch::gfx906),
2758  key_type,
2759  value_type,
2760  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2761  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
2762  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2763  : radix_sort_onesweep_config<kernel_config<512, 12>,
2764  kernel_config<512, 12>,
2765  8,
2766  block_radix_rank_algorithm::match>
2767 {};
2768 
2769 // Based on key_type = float, value_type = empty_type
2770 template<class key_type, class value_type>
2771 struct default_radix_sort_onesweep_config<
2772  static_cast<unsigned int>(target_arch::gfx906),
2773  key_type,
2774  value_type,
2775  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2776  && (sizeof(key_type) > 2)
2777  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2778  : radix_sort_onesweep_config<kernel_config<512, 12>,
2779  kernel_config<512, 12>,
2780  8,
2781  block_radix_rank_algorithm::match>
2782 {};
2783 
2784 // Based on key_type = rocprim::half, value_type = int64_t
2785 template<class key_type, class value_type>
2786 struct default_radix_sort_onesweep_config<
2787  static_cast<unsigned int>(target_arch::gfx906),
2788  key_type,
2789  value_type,
2790  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2791  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
2792  : radix_sort_onesweep_config<kernel_config<512, 6>,
2793  kernel_config<512, 6>,
2794  8,
2795  block_radix_rank_algorithm::match>
2796 {};
2797 
2798 // Based on key_type = rocprim::half, value_type = int
2799 template<class key_type, class value_type>
2800 struct default_radix_sort_onesweep_config<
2801  static_cast<unsigned int>(target_arch::gfx906),
2802  key_type,
2803  value_type,
2804  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2805  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
2806  : radix_sort_onesweep_config<kernel_config<256, 16>,
2807  kernel_config<256, 16>,
2808  8,
2809  block_radix_rank_algorithm::match>
2810 {};
2811 
2812 // Based on key_type = rocprim::half, value_type = short
2813 template<class key_type, class value_type>
2814 struct default_radix_sort_onesweep_config<
2815  static_cast<unsigned int>(target_arch::gfx906),
2816  key_type,
2817  value_type,
2818  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2819  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
2820  : radix_sort_onesweep_config<kernel_config<256, 16>,
2821  kernel_config<256, 16>,
2822  8,
2823  block_radix_rank_algorithm::match>
2824 {};
2825 
2826 // Based on key_type = rocprim::half, value_type = int8_t
2827 template<class key_type, class value_type>
2828 struct default_radix_sort_onesweep_config<
2829  static_cast<unsigned int>(target_arch::gfx906),
2830  key_type,
2831  value_type,
2832  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2833  && (sizeof(value_type) <= 1)
2834  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2835  : radix_sort_onesweep_config<kernel_config<512, 12>,
2836  kernel_config<512, 12>,
2837  8,
2838  block_radix_rank_algorithm::match>
2839 {};
2840 
2841 // Based on key_type = rocprim::half, value_type = empty_type
2842 template<class key_type, class value_type>
2843 struct default_radix_sort_onesweep_config<
2844  static_cast<unsigned int>(target_arch::gfx906),
2845  key_type,
2846  value_type,
2847  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2848  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2849  : radix_sort_onesweep_config<kernel_config<512, 12>,
2850  kernel_config<512, 12>,
2851  8,
2852  block_radix_rank_algorithm::match>
2853 {};
2854 
2855 // Based on key_type = int64_t, value_type = int64_t
2856 template<class key_type, class value_type>
2857 struct default_radix_sort_onesweep_config<
2858  static_cast<unsigned int>(target_arch::gfx906),
2859  key_type,
2860  value_type,
2861  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2862  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
2863  && (sizeof(value_type) > 4))>>
2864  : radix_sort_onesweep_config<kernel_config<512, 6>,
2865  kernel_config<512, 6>,
2866  8,
2867  block_radix_rank_algorithm::match>
2868 {};
2869 
2870 // Based on key_type = int64_t, value_type = int
2871 template<class key_type, class value_type>
2872 struct default_radix_sort_onesweep_config<
2873  static_cast<unsigned int>(target_arch::gfx906),
2874  key_type,
2875  value_type,
2876  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2877  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
2878  && (sizeof(value_type) > 2))>>
2879  : radix_sort_onesweep_config<kernel_config<256, 12>,
2880  kernel_config<256, 12>,
2881  8,
2882  block_radix_rank_algorithm::match>
2883 {};
2884 
2885 // Based on key_type = int64_t, value_type = short
2886 template<class key_type, class value_type>
2887 struct default_radix_sort_onesweep_config<
2888  static_cast<unsigned int>(target_arch::gfx906),
2889  key_type,
2890  value_type,
2891  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2892  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
2893  && (sizeof(value_type) > 1))>>
2894  : radix_sort_onesweep_config<kernel_config<512, 6>,
2895  kernel_config<512, 6>,
2896  8,
2897  block_radix_rank_algorithm::match>
2898 {};
2899 
2900 // Based on key_type = int64_t, value_type = int8_t
2901 template<class key_type, class value_type>
2902 struct default_radix_sort_onesweep_config<
2903  static_cast<unsigned int>(target_arch::gfx906),
2904  key_type,
2905  value_type,
2906  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2907  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
2908  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2909  : radix_sort_onesweep_config<kernel_config<512, 6>,
2910  kernel_config<512, 6>,
2911  8,
2912  block_radix_rank_algorithm::match>
2913 {};
2914 
2915 // Based on key_type = int64_t, value_type = empty_type
2916 template<class key_type, class value_type>
2917 struct default_radix_sort_onesweep_config<
2918  static_cast<unsigned int>(target_arch::gfx906),
2919  key_type,
2920  value_type,
2921  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2922  && (sizeof(key_type) > 4)
2923  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2924  : radix_sort_onesweep_config<kernel_config<256, 12>,
2925  kernel_config<256, 12>,
2926  8,
2927  block_radix_rank_algorithm::match>
2928 {};
2929 
2930 // Based on key_type = int, value_type = int64_t
2931 template<class key_type, class value_type>
2932 struct default_radix_sort_onesweep_config<
2933  static_cast<unsigned int>(target_arch::gfx906),
2934  key_type,
2935  value_type,
2936  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2937  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
2938  && (sizeof(value_type) > 4))>>
2939  : radix_sort_onesweep_config<kernel_config<256, 12>,
2940  kernel_config<256, 12>,
2941  8,
2942  block_radix_rank_algorithm::match>
2943 {};
2944 
2945 // Based on key_type = int, value_type = int
2946 template<class key_type, class value_type>
2947 struct default_radix_sort_onesweep_config<
2948  static_cast<unsigned int>(target_arch::gfx906),
2949  key_type,
2950  value_type,
2951  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2952  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
2953  && (sizeof(value_type) > 2))>>
2954  : radix_sort_onesweep_config<kernel_config<512, 12>,
2955  kernel_config<512, 12>,
2956  5,
2957  block_radix_rank_algorithm::match>
2958 {};
2959 
2960 // Based on key_type = int, value_type = short
2961 template<class key_type, class value_type>
2962 struct default_radix_sort_onesweep_config<
2963  static_cast<unsigned int>(target_arch::gfx906),
2964  key_type,
2965  value_type,
2966  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2967  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
2968  && (sizeof(value_type) > 1))>>
2969  : radix_sort_onesweep_config<kernel_config<256, 18>,
2970  kernel_config<256, 18>,
2971  8,
2972  block_radix_rank_algorithm::match>
2973 {};
2974 
2975 // Based on key_type = int, value_type = int8_t
2976 template<class key_type, class value_type>
2977 struct default_radix_sort_onesweep_config<
2978  static_cast<unsigned int>(target_arch::gfx906),
2979  key_type,
2980  value_type,
2981  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2982  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
2983  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
2984  : radix_sort_onesweep_config<kernel_config<256, 18>,
2985  kernel_config<256, 18>,
2986  8,
2987  block_radix_rank_algorithm::match>
2988 {};
2989 
2990 // Based on key_type = int, value_type = empty_type
2991 template<class key_type, class value_type>
2992 struct default_radix_sort_onesweep_config<
2993  static_cast<unsigned int>(target_arch::gfx906),
2994  key_type,
2995  value_type,
2996  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2997  && (sizeof(key_type) > 2)
2998  && (std::is_same<value_type, rocprim::empty_type>::value))>>
2999  : radix_sort_onesweep_config<kernel_config<256, 22>,
3000  kernel_config<256, 22>,
3001  8,
3002  block_radix_rank_algorithm::match>
3003 {};
3004 
3005 // Based on key_type = short, value_type = int64_t
3006 template<class key_type, class value_type>
3007 struct default_radix_sort_onesweep_config<
3008  static_cast<unsigned int>(target_arch::gfx906),
3009  key_type,
3010  value_type,
3011  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
3012  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
3013  && (sizeof(value_type) > 4))>>
3014  : radix_sort_onesweep_config<kernel_config<256, 12>,
3015  kernel_config<256, 12>,
3016  8,
3017  block_radix_rank_algorithm::match>
3018 {};
3019 
3020 // Based on key_type = short, value_type = int
3021 template<class key_type, class value_type>
3022 struct default_radix_sort_onesweep_config<
3023  static_cast<unsigned int>(target_arch::gfx906),
3024  key_type,
3025  value_type,
3026  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
3027  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
3028  && (sizeof(value_type) > 2))>>
3029  : radix_sort_onesweep_config<kernel_config<512, 22>,
3030  kernel_config<512, 22>,
3031  8,
3032  block_radix_rank_algorithm::match>
3033 {};
3034 
3035 // Based on key_type = short, value_type = short
3036 template<class key_type, class value_type>
3037 struct default_radix_sort_onesweep_config<
3038  static_cast<unsigned int>(target_arch::gfx906),
3039  key_type,
3040  value_type,
3041  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
3042  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
3043  && (sizeof(value_type) > 1))>>
3044  : radix_sort_onesweep_config<kernel_config<512, 12>,
3045  kernel_config<512, 12>,
3046  4,
3047  block_radix_rank_algorithm::match>
3048 {};
3049 
3050 // Based on key_type = short, value_type = int8_t
3051 template<class key_type, class value_type>
3052 struct default_radix_sort_onesweep_config<
3053  static_cast<unsigned int>(target_arch::gfx906),
3054  key_type,
3055  value_type,
3056  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
3057  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
3058  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
3059  : radix_sort_onesweep_config<kernel_config<512, 12>,
3060  kernel_config<512, 12>,
3061  8,
3062  block_radix_rank_algorithm::match>
3063 {};
3064 
3065 // Based on key_type = short, value_type = empty_type
3066 template<class key_type, class value_type>
3067 struct default_radix_sort_onesweep_config<
3068  static_cast<unsigned int>(target_arch::gfx906),
3069  key_type,
3070  value_type,
3071  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
3072  && (sizeof(key_type) > 1)
3073  && (std::is_same<value_type, rocprim::empty_type>::value))>>
3074  : radix_sort_onesweep_config<kernel_config<512, 12>,
3075  kernel_config<512, 12>,
3076  8,
3077  block_radix_rank_algorithm::match>
3078 {};
3079 
3080 // Based on key_type = int8_t, value_type = int64_t
3081 template<class key_type, class value_type>
3082 struct default_radix_sort_onesweep_config<
3083  static_cast<unsigned int>(target_arch::gfx906),
3084  key_type,
3085  value_type,
3086  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
3087  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
3088  : radix_sort_onesweep_config<kernel_config<512, 6>,
3089  kernel_config<512, 6>,
3090  8,
3091  block_radix_rank_algorithm::match>
3092 {};
3093 
3094 // Based on key_type = int8_t, value_type = int
3095 template<class key_type, class value_type>
3096 struct default_radix_sort_onesweep_config<
3097  static_cast<unsigned int>(target_arch::gfx906),
3098  key_type,
3099  value_type,
3100  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
3101  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
3102  : radix_sort_onesweep_config<kernel_config<256, 16>,
3103  kernel_config<256, 16>,
3104  8,
3105  block_radix_rank_algorithm::match>
3106 {};
3107 
3108 // Based on key_type = int8_t, value_type = short
3109 template<class key_type, class value_type>
3110 struct default_radix_sort_onesweep_config<
3111  static_cast<unsigned int>(target_arch::gfx906),
3112  key_type,
3113  value_type,
3114  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
3115  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
3116  : radix_sort_onesweep_config<kernel_config<512, 8>,
3117  kernel_config<512, 8>,
3118  8,
3119  block_radix_rank_algorithm::match>
3120 {};
3121 
3122 // Based on key_type = int8_t, value_type = int8_t
3123 template<class key_type, class value_type>
3124 struct default_radix_sort_onesweep_config<
3125  static_cast<unsigned int>(target_arch::gfx906),
3126  key_type,
3127  value_type,
3128  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
3129  && (sizeof(value_type) <= 1)
3130  && (!std::is_same<value_type, rocprim::empty_type>::value))>>
3131  : radix_sort_onesweep_config<kernel_config<512, 16>,
3132  kernel_config<512, 16>,
3133  4,
3134  block_radix_rank_algorithm::basic>
3135 {};
3136 
3137 // Based on key_type = int8_t, value_type = empty_type
3138 template<class key_type, class value_type>
3139 struct default_radix_sort_onesweep_config<
3140  static_cast<unsigned int>(target_arch::gfx906),
3141  key_type,
3142  value_type,
3143  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
3144  && (std::is_same<value_type, rocprim::empty_type>::value))>>
3145  : radix_sort_onesweep_config<kernel_config<512, 12>,
3146  kernel_config<512, 12>,
3147  8,
3148  block_radix_rank_algorithm::match>
3149 {};
3150 
3151 } // end namespace detail
3152 
3153 END_ROCPRIM_NAMESPACE
3154 
3156 // end of group primitivesmodule_deviceconfigs
3157 
3158 #endif // ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_RADIX_SORT_ONESWEEP_HPP_
detail::default_radix_sort_onesweep_config
Definition: device_radix_sort_onesweep.hpp:45
detail
Deprecated: Configuration of device-level scan primitives.
Definition: block_histogram.hpp:62
radix_sort_onesweep_config
Configuration of subalgorithm Onesweep.
Definition: device_config_helper.hpp:188
detail::radix_sort_onesweep_config_base
Definition: device_config_helper.hpp:211
Generated by   doxygen 1.8.13