rocPRIM
device_scan_by_key.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_SCAN_BY_KEY_HPP_
22 #define ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_SCAN_BY_KEY_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, class key_type, class value_type, class enable = void>
42 struct default_scan_by_key_config : default_scan_by_key_config_base<key_type, value_type>
43 {};
44 
45 // Based on key_type = double, value_type = int64_t
46 template<class key_type, class value_type>
47 struct default_scan_by_key_config<
48  static_cast<unsigned int>(target_arch::gfx908),
49  key_type,
50  value_type,
51  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
52  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
53  && (sizeof(value_type) > 4))>>
54  : scan_by_key_config_v2<256,
55  14,
56  ::rocprim::block_load_method::block_load_transpose,
57  ::rocprim::block_store_method::block_store_transpose,
58  block_scan_algorithm::using_warp_scan>
59 {};
60 
61 // Based on key_type = double, value_type = int
62 template<class key_type, class value_type>
63 struct default_scan_by_key_config<
64  static_cast<unsigned int>(target_arch::gfx908),
65  key_type,
66  value_type,
67  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
68  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
69  && (sizeof(value_type) > 2))>>
70  : scan_by_key_config_v2<256,
71  14,
72  ::rocprim::block_load_method::block_load_transpose,
73  ::rocprim::block_store_method::block_store_transpose,
74  block_scan_algorithm::reduce_then_scan>
75 {};
76 
77 // Based on key_type = double, value_type = short
78 template<class key_type, class value_type>
79 struct default_scan_by_key_config<
80  static_cast<unsigned int>(target_arch::gfx908),
81  key_type,
82  value_type,
83  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
84  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
85  && (sizeof(value_type) > 1))>>
86  : scan_by_key_config_v2<64,
87  14,
88  ::rocprim::block_load_method::block_load_transpose,
89  ::rocprim::block_store_method::block_store_transpose,
90  block_scan_algorithm::using_warp_scan>
91 {};
92 
93 // Based on key_type = double, value_type = int8_t
94 template<class key_type, class value_type>
95 struct default_scan_by_key_config<
96  static_cast<unsigned int>(target_arch::gfx908),
97  key_type,
98  value_type,
99  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
100  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
101  : scan_by_key_config_v2<64,
102  14,
103  ::rocprim::block_load_method::block_load_transpose,
104  ::rocprim::block_store_method::block_store_transpose,
105  block_scan_algorithm::reduce_then_scan>
106 {};
107 
108 // Based on key_type = float, value_type = int64_t
109 template<class key_type, class value_type>
110 struct default_scan_by_key_config<
111  static_cast<unsigned int>(target_arch::gfx908),
112  key_type,
113  value_type,
114  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
115  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
116  && (sizeof(value_type) > 4))>>
117  : scan_by_key_config_v2<128,
118  10,
119  ::rocprim::block_load_method::block_load_transpose,
120  ::rocprim::block_store_method::block_store_transpose,
121  block_scan_algorithm::using_warp_scan>
122 {};
123 
124 // Based on key_type = float, value_type = int
125 template<class key_type, class value_type>
126 struct default_scan_by_key_config<
127  static_cast<unsigned int>(target_arch::gfx908),
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) <= 4)
132  && (sizeof(value_type) > 2))>>
133  : scan_by_key_config_v2<256,
134  12,
135  ::rocprim::block_load_method::block_load_transpose,
136  ::rocprim::block_store_method::block_store_transpose,
137  block_scan_algorithm::reduce_then_scan>
138 {};
139 
140 // Based on key_type = float, value_type = short
141 template<class key_type, class value_type>
142 struct default_scan_by_key_config<
143  static_cast<unsigned int>(target_arch::gfx908),
144  key_type,
145  value_type,
146  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
147  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
148  && (sizeof(value_type) > 1))>>
149  : scan_by_key_config_v2<256,
150  14,
151  ::rocprim::block_load_method::block_load_transpose,
152  ::rocprim::block_store_method::block_store_transpose,
153  block_scan_algorithm::reduce_then_scan>
154 {};
155 
156 // Based on key_type = float, value_type = int8_t
157 template<class key_type, class value_type>
158 struct default_scan_by_key_config<
159  static_cast<unsigned int>(target_arch::gfx908),
160  key_type,
161  value_type,
162  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
163  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
164  : scan_by_key_config_v2<128,
165  14,
166  ::rocprim::block_load_method::block_load_transpose,
167  ::rocprim::block_store_method::block_store_transpose,
168  block_scan_algorithm::using_warp_scan>
169 {};
170 
171 // Based on key_type = rocprim::half, value_type = int64_t
172 template<class key_type, class value_type>
173 struct default_scan_by_key_config<
174  static_cast<unsigned int>(target_arch::gfx908),
175  key_type,
176  value_type,
177  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
178  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
179  : scan_by_key_config_v2<64,
180  18,
181  ::rocprim::block_load_method::block_load_transpose,
182  ::rocprim::block_store_method::block_store_transpose,
183  block_scan_algorithm::using_warp_scan>
184 {};
185 
186 // Based on key_type = rocprim::half, value_type = int
187 template<class key_type, class value_type>
188 struct default_scan_by_key_config<
189  static_cast<unsigned int>(target_arch::gfx908),
190  key_type,
191  value_type,
192  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
193  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
194  : scan_by_key_config_v2<256,
195  18,
196  ::rocprim::block_load_method::block_load_transpose,
197  ::rocprim::block_store_method::block_store_transpose,
198  block_scan_algorithm::using_warp_scan>
199 {};
200 
201 // Based on key_type = rocprim::half, value_type = short
202 template<class key_type, class value_type>
203 struct default_scan_by_key_config<
204  static_cast<unsigned int>(target_arch::gfx908),
205  key_type,
206  value_type,
207  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
208  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
209  : scan_by_key_config_v2<256,
210  20,
211  ::rocprim::block_load_method::block_load_transpose,
212  ::rocprim::block_store_method::block_store_transpose,
213  block_scan_algorithm::reduce_then_scan>
214 {};
215 
216 // Based on key_type = rocprim::half, value_type = int8_t
217 template<class key_type, class value_type>
218 struct default_scan_by_key_config<
219  static_cast<unsigned int>(target_arch::gfx908),
220  key_type,
221  value_type,
222  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
223  && (sizeof(value_type) <= 1))>>
224  : scan_by_key_config_v2<256,
225  24,
226  ::rocprim::block_load_method::block_load_transpose,
227  ::rocprim::block_store_method::block_store_transpose,
228  block_scan_algorithm::using_warp_scan>
229 {};
230 
231 // Based on key_type = int64_t, value_type = int64_t
232 template<class key_type, class value_type>
233 struct default_scan_by_key_config<
234  static_cast<unsigned int>(target_arch::gfx908),
235  key_type,
236  value_type,
237  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
238  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
239  && (sizeof(value_type) > 4))>>
240  : scan_by_key_config_v2<256,
241  14,
242  ::rocprim::block_load_method::block_load_transpose,
243  ::rocprim::block_store_method::block_store_transpose,
244  block_scan_algorithm::reduce_then_scan>
245 {};
246 
247 // Based on key_type = int64_t, value_type = int
248 template<class key_type, class value_type>
249 struct default_scan_by_key_config<
250  static_cast<unsigned int>(target_arch::gfx908),
251  key_type,
252  value_type,
253  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
254  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
255  && (sizeof(value_type) > 2))>>
256  : scan_by_key_config_v2<256,
257  14,
258  ::rocprim::block_load_method::block_load_transpose,
259  ::rocprim::block_store_method::block_store_transpose,
260  block_scan_algorithm::reduce_then_scan>
261 {};
262 
263 // Based on key_type = int64_t, value_type = short
264 template<class key_type, class value_type>
265 struct default_scan_by_key_config<
266  static_cast<unsigned int>(target_arch::gfx908),
267  key_type,
268  value_type,
269  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
270  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
271  && (sizeof(value_type) > 1))>>
272  : scan_by_key_config_v2<64,
273  14,
274  ::rocprim::block_load_method::block_load_transpose,
275  ::rocprim::block_store_method::block_store_transpose,
276  block_scan_algorithm::using_warp_scan>
277 {};
278 
279 // Based on key_type = int64_t, value_type = int8_t
280 template<class key_type, class value_type>
281 struct default_scan_by_key_config<
282  static_cast<unsigned int>(target_arch::gfx908),
283  key_type,
284  value_type,
285  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
286  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
287  : scan_by_key_config_v2<256,
288  14,
289  ::rocprim::block_load_method::block_load_transpose,
290  ::rocprim::block_store_method::block_store_transpose,
291  block_scan_algorithm::using_warp_scan>
292 {};
293 
294 // Based on key_type = int, value_type = int64_t
295 template<class key_type, class value_type>
296 struct default_scan_by_key_config<
297  static_cast<unsigned int>(target_arch::gfx908),
298  key_type,
299  value_type,
300  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
301  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
302  && (sizeof(value_type) > 4))>>
303  : scan_by_key_config_v2<128,
304  10,
305  ::rocprim::block_load_method::block_load_transpose,
306  ::rocprim::block_store_method::block_store_transpose,
307  block_scan_algorithm::using_warp_scan>
308 {};
309 
310 // Based on key_type = int, value_type = int
311 template<class key_type, class value_type>
312 struct default_scan_by_key_config<
313  static_cast<unsigned int>(target_arch::gfx908),
314  key_type,
315  value_type,
316  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
317  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
318  && (sizeof(value_type) > 2))>>
319  : scan_by_key_config_v2<256,
320  12,
321  ::rocprim::block_load_method::block_load_transpose,
322  ::rocprim::block_store_method::block_store_transpose,
323  block_scan_algorithm::reduce_then_scan>
324 {};
325 
326 // Based on key_type = int, value_type = short
327 template<class key_type, class value_type>
328 struct default_scan_by_key_config<
329  static_cast<unsigned int>(target_arch::gfx908),
330  key_type,
331  value_type,
332  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
333  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
334  && (sizeof(value_type) > 1))>>
335  : scan_by_key_config_v2<256,
336  14,
337  ::rocprim::block_load_method::block_load_transpose,
338  ::rocprim::block_store_method::block_store_transpose,
339  block_scan_algorithm::reduce_then_scan>
340 {};
341 
342 // Based on key_type = int, value_type = int8_t
343 template<class key_type, class value_type>
344 struct default_scan_by_key_config<
345  static_cast<unsigned int>(target_arch::gfx908),
346  key_type,
347  value_type,
348  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
349  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
350  : scan_by_key_config_v2<128,
351  14,
352  ::rocprim::block_load_method::block_load_transpose,
353  ::rocprim::block_store_method::block_store_transpose,
354  block_scan_algorithm::using_warp_scan>
355 {};
356 
357 // Based on key_type = short, value_type = int64_t
358 template<class key_type, class value_type>
359 struct default_scan_by_key_config<
360  static_cast<unsigned int>(target_arch::gfx908),
361  key_type,
362  value_type,
363  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
364  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
365  && (sizeof(value_type) > 4))>>
366  : scan_by_key_config_v2<64,
367  18,
368  ::rocprim::block_load_method::block_load_transpose,
369  ::rocprim::block_store_method::block_store_transpose,
370  block_scan_algorithm::reduce_then_scan>
371 {};
372 
373 // Based on key_type = short, value_type = int
374 template<class key_type, class value_type>
375 struct default_scan_by_key_config<
376  static_cast<unsigned int>(target_arch::gfx908),
377  key_type,
378  value_type,
379  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
380  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
381  && (sizeof(value_type) > 2))>>
382  : scan_by_key_config_v2<256,
383  18,
384  ::rocprim::block_load_method::block_load_transpose,
385  ::rocprim::block_store_method::block_store_transpose,
386  block_scan_algorithm::using_warp_scan>
387 {};
388 
389 // Based on key_type = short, value_type = short
390 template<class key_type, class value_type>
391 struct default_scan_by_key_config<
392  static_cast<unsigned int>(target_arch::gfx908),
393  key_type,
394  value_type,
395  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
396  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
397  && (sizeof(value_type) > 1))>>
398  : scan_by_key_config_v2<256,
399  20,
400  ::rocprim::block_load_method::block_load_transpose,
401  ::rocprim::block_store_method::block_store_transpose,
402  block_scan_algorithm::reduce_then_scan>
403 {};
404 
405 // Based on key_type = short, value_type = int8_t
406 template<class key_type, class value_type>
407 struct default_scan_by_key_config<
408  static_cast<unsigned int>(target_arch::gfx908),
409  key_type,
410  value_type,
411  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
412  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1))>>
413  : scan_by_key_config_v2<256,
414  24,
415  ::rocprim::block_load_method::block_load_transpose,
416  ::rocprim::block_store_method::block_store_transpose,
417  block_scan_algorithm::using_warp_scan>
418 {};
419 
420 // Based on key_type = int8_t, value_type = int64_t
421 template<class key_type, class value_type>
422 struct default_scan_by_key_config<
423  static_cast<unsigned int>(target_arch::gfx908),
424  key_type,
425  value_type,
426  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
427  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
428  : scan_by_key_config_v2<64,
429  18,
430  ::rocprim::block_load_method::block_load_transpose,
431  ::rocprim::block_store_method::block_store_transpose,
432  block_scan_algorithm::using_warp_scan>
433 {};
434 
435 // Based on key_type = int8_t, value_type = int
436 template<class key_type, class value_type>
437 struct default_scan_by_key_config<
438  static_cast<unsigned int>(target_arch::gfx908),
439  key_type,
440  value_type,
441  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
442  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
443  : scan_by_key_config_v2<128,
444  18,
445  ::rocprim::block_load_method::block_load_transpose,
446  ::rocprim::block_store_method::block_store_transpose,
447  block_scan_algorithm::using_warp_scan>
448 {};
449 
450 // Based on key_type = int8_t, value_type = short
451 template<class key_type, class value_type>
452 struct default_scan_by_key_config<
453  static_cast<unsigned int>(target_arch::gfx908),
454  key_type,
455  value_type,
456  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
457  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
458  : scan_by_key_config_v2<256,
459  20,
460  ::rocprim::block_load_method::block_load_transpose,
461  ::rocprim::block_store_method::block_store_transpose,
462  block_scan_algorithm::reduce_then_scan>
463 {};
464 
465 // Based on key_type = int8_t, value_type = int8_t
466 template<class key_type, class value_type>
467 struct default_scan_by_key_config<
468  static_cast<unsigned int>(target_arch::gfx908),
469  key_type,
470  value_type,
471  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
472  && (sizeof(value_type) <= 1))>>
473  : scan_by_key_config_v2<256,
474  24,
475  ::rocprim::block_load_method::block_load_transpose,
476  ::rocprim::block_store_method::block_store_transpose,
477  block_scan_algorithm::reduce_then_scan>
478 {};
479 
480 // Based on key_type = double, value_type = int64_t
481 template<class key_type, class value_type>
482 struct default_scan_by_key_config<
483  static_cast<unsigned int>(target_arch::gfx900),
484  key_type,
485  value_type,
486  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
487  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
488  && (sizeof(value_type) > 4))>>
489  : scan_by_key_config_v2<256,
490  14,
491  ::rocprim::block_load_method::block_load_transpose,
492  ::rocprim::block_store_method::block_store_transpose,
493  block_scan_algorithm::reduce_then_scan>
494 {};
495 
496 // Based on key_type = double, value_type = int
497 template<class key_type, class value_type>
498 struct default_scan_by_key_config<
499  static_cast<unsigned int>(target_arch::gfx900),
500  key_type,
501  value_type,
502  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
503  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
504  && (sizeof(value_type) > 2))>>
505  : scan_by_key_config_v2<128,
506  19,
507  ::rocprim::block_load_method::block_load_transpose,
508  ::rocprim::block_store_method::block_store_transpose,
509  block_scan_algorithm::reduce_then_scan>
510 {};
511 
512 // Based on key_type = double, value_type = short
513 template<class key_type, class value_type>
514 struct default_scan_by_key_config<
515  static_cast<unsigned int>(target_arch::gfx900),
516  key_type,
517  value_type,
518  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
519  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
520  && (sizeof(value_type) > 1))>>
521  : scan_by_key_config_v2<128,
522  14,
523  ::rocprim::block_load_method::block_load_transpose,
524  ::rocprim::block_store_method::block_store_transpose,
525  block_scan_algorithm::reduce_then_scan>
526 {};
527 
528 // Based on key_type = double, value_type = int8_t
529 template<class key_type, class value_type>
530 struct default_scan_by_key_config<
531  static_cast<unsigned int>(target_arch::gfx900),
532  key_type,
533  value_type,
534  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
535  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
536  : scan_by_key_config_v2<128,
537  14,
538  ::rocprim::block_load_method::block_load_transpose,
539  ::rocprim::block_store_method::block_store_transpose,
540  block_scan_algorithm::using_warp_scan>
541 {};
542 
543 // Based on key_type = float, value_type = int64_t
544 template<class key_type, class value_type>
545 struct default_scan_by_key_config<
546  static_cast<unsigned int>(target_arch::gfx900),
547  key_type,
548  value_type,
549  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
550  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
551  && (sizeof(value_type) > 4))>>
552  : scan_by_key_config_v2<256,
553  10,
554  ::rocprim::block_load_method::block_load_transpose,
555  ::rocprim::block_store_method::block_store_transpose,
556  block_scan_algorithm::reduce_then_scan>
557 {};
558 
559 // Based on key_type = float, value_type = int
560 template<class key_type, class value_type>
561 struct default_scan_by_key_config<
562  static_cast<unsigned int>(target_arch::gfx900),
563  key_type,
564  value_type,
565  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
566  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
567  && (sizeof(value_type) > 2))>>
568  : scan_by_key_config_v2<256,
569  12,
570  ::rocprim::block_load_method::block_load_transpose,
571  ::rocprim::block_store_method::block_store_transpose,
572  block_scan_algorithm::reduce_then_scan>
573 {};
574 
575 // Based on key_type = float, value_type = short
576 template<class key_type, class value_type>
577 struct default_scan_by_key_config<
578  static_cast<unsigned int>(target_arch::gfx900),
579  key_type,
580  value_type,
581  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
582  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
583  && (sizeof(value_type) > 1))>>
584  : scan_by_key_config_v2<256,
585  14,
586  ::rocprim::block_load_method::block_load_transpose,
587  ::rocprim::block_store_method::block_store_transpose,
588  block_scan_algorithm::reduce_then_scan>
589 {};
590 
591 // Based on key_type = float, value_type = int8_t
592 template<class key_type, class value_type>
593 struct default_scan_by_key_config<
594  static_cast<unsigned int>(target_arch::gfx900),
595  key_type,
596  value_type,
597  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
598  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
599  : scan_by_key_config_v2<256,
600  14,
601  ::rocprim::block_load_method::block_load_transpose,
602  ::rocprim::block_store_method::block_store_transpose,
603  block_scan_algorithm::using_warp_scan>
604 {};
605 
606 // Based on key_type = rocprim::half, value_type = int64_t
607 template<class key_type, class value_type>
608 struct default_scan_by_key_config<
609  static_cast<unsigned int>(target_arch::gfx900),
610  key_type,
611  value_type,
612  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
613  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
614  : scan_by_key_config_v2<64,
615  18,
616  ::rocprim::block_load_method::block_load_transpose,
617  ::rocprim::block_store_method::block_store_transpose,
618  block_scan_algorithm::reduce_then_scan>
619 {};
620 
621 // Based on key_type = rocprim::half, value_type = int
622 template<class key_type, class value_type>
623 struct default_scan_by_key_config<
624  static_cast<unsigned int>(target_arch::gfx900),
625  key_type,
626  value_type,
627  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
628  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
629  : scan_by_key_config_v2<256,
630  18,
631  ::rocprim::block_load_method::block_load_transpose,
632  ::rocprim::block_store_method::block_store_transpose,
633  block_scan_algorithm::using_warp_scan>
634 {};
635 
636 // Based on key_type = rocprim::half, value_type = short
637 template<class key_type, class value_type>
638 struct default_scan_by_key_config<
639  static_cast<unsigned int>(target_arch::gfx900),
640  key_type,
641  value_type,
642  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
643  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
644  : scan_by_key_config_v2<256,
645  22,
646  ::rocprim::block_load_method::block_load_transpose,
647  ::rocprim::block_store_method::block_store_transpose,
648  block_scan_algorithm::reduce_then_scan>
649 {};
650 
651 // Based on key_type = rocprim::half, value_type = int8_t
652 template<class key_type, class value_type>
653 struct default_scan_by_key_config<
654  static_cast<unsigned int>(target_arch::gfx900),
655  key_type,
656  value_type,
657  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
658  && (sizeof(value_type) <= 1))>>
659  : scan_by_key_config_v2<256,
660  20,
661  ::rocprim::block_load_method::block_load_transpose,
662  ::rocprim::block_store_method::block_store_transpose,
663  block_scan_algorithm::using_warp_scan>
664 {};
665 
666 // Based on key_type = int64_t, value_type = int64_t
667 template<class key_type, class value_type>
668 struct default_scan_by_key_config<
669  static_cast<unsigned int>(target_arch::gfx900),
670  key_type,
671  value_type,
672  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
673  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
674  && (sizeof(value_type) > 4))>>
675  : scan_by_key_config_v2<256,
676  14,
677  ::rocprim::block_load_method::block_load_transpose,
678  ::rocprim::block_store_method::block_store_transpose,
679  block_scan_algorithm::using_warp_scan>
680 {};
681 
682 // Based on key_type = int64_t, value_type = int
683 template<class key_type, class value_type>
684 struct default_scan_by_key_config<
685  static_cast<unsigned int>(target_arch::gfx900),
686  key_type,
687  value_type,
688  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
689  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
690  && (sizeof(value_type) > 2))>>
691  : scan_by_key_config_v2<128,
692  19,
693  ::rocprim::block_load_method::block_load_transpose,
694  ::rocprim::block_store_method::block_store_transpose,
695  block_scan_algorithm::reduce_then_scan>
696 {};
697 
698 // Based on key_type = int64_t, value_type = short
699 template<class key_type, class value_type>
700 struct default_scan_by_key_config<
701  static_cast<unsigned int>(target_arch::gfx900),
702  key_type,
703  value_type,
704  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
705  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
706  && (sizeof(value_type) > 1))>>
707  : scan_by_key_config_v2<128,
708  14,
709  ::rocprim::block_load_method::block_load_transpose,
710  ::rocprim::block_store_method::block_store_transpose,
711  block_scan_algorithm::reduce_then_scan>
712 {};
713 
714 // Based on key_type = int64_t, value_type = int8_t
715 template<class key_type, class value_type>
716 struct default_scan_by_key_config<
717  static_cast<unsigned int>(target_arch::gfx900),
718  key_type,
719  value_type,
720  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
721  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
722  : scan_by_key_config_v2<128,
723  14,
724  ::rocprim::block_load_method::block_load_transpose,
725  ::rocprim::block_store_method::block_store_transpose,
726  block_scan_algorithm::using_warp_scan>
727 {};
728 
729 // Based on key_type = int, value_type = int64_t
730 template<class key_type, class value_type>
731 struct default_scan_by_key_config<
732  static_cast<unsigned int>(target_arch::gfx900),
733  key_type,
734  value_type,
735  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
736  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
737  && (sizeof(value_type) > 4))>>
738  : scan_by_key_config_v2<256,
739  10,
740  ::rocprim::block_load_method::block_load_transpose,
741  ::rocprim::block_store_method::block_store_transpose,
742  block_scan_algorithm::using_warp_scan>
743 {};
744 
745 // Based on key_type = int, value_type = int
746 template<class key_type, class value_type>
747 struct default_scan_by_key_config<
748  static_cast<unsigned int>(target_arch::gfx900),
749  key_type,
750  value_type,
751  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
752  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
753  && (sizeof(value_type) > 2))>>
754  : scan_by_key_config_v2<256,
755  12,
756  ::rocprim::block_load_method::block_load_transpose,
757  ::rocprim::block_store_method::block_store_transpose,
758  block_scan_algorithm::reduce_then_scan>
759 {};
760 
761 // Based on key_type = int, value_type = short
762 template<class key_type, class value_type>
763 struct default_scan_by_key_config<
764  static_cast<unsigned int>(target_arch::gfx900),
765  key_type,
766  value_type,
767  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
768  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
769  && (sizeof(value_type) > 1))>>
770  : scan_by_key_config_v2<256,
771  14,
772  ::rocprim::block_load_method::block_load_transpose,
773  ::rocprim::block_store_method::block_store_transpose,
774  block_scan_algorithm::reduce_then_scan>
775 {};
776 
777 // Based on key_type = int, value_type = int8_t
778 template<class key_type, class value_type>
779 struct default_scan_by_key_config<
780  static_cast<unsigned int>(target_arch::gfx900),
781  key_type,
782  value_type,
783  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
784  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
785  : scan_by_key_config_v2<256,
786  14,
787  ::rocprim::block_load_method::block_load_transpose,
788  ::rocprim::block_store_method::block_store_transpose,
789  block_scan_algorithm::using_warp_scan>
790 {};
791 
792 // Based on key_type = short, value_type = int64_t
793 template<class key_type, class value_type>
794 struct default_scan_by_key_config<
795  static_cast<unsigned int>(target_arch::gfx900),
796  key_type,
797  value_type,
798  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
799  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
800  && (sizeof(value_type) > 4))>>
801  : scan_by_key_config_v2<64,
802  18,
803  ::rocprim::block_load_method::block_load_transpose,
804  ::rocprim::block_store_method::block_store_transpose,
805  block_scan_algorithm::using_warp_scan>
806 {};
807 
808 // Based on key_type = short, value_type = int
809 template<class key_type, class value_type>
810 struct default_scan_by_key_config<
811  static_cast<unsigned int>(target_arch::gfx900),
812  key_type,
813  value_type,
814  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
815  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
816  && (sizeof(value_type) > 2))>>
817  : scan_by_key_config_v2<256,
818  18,
819  ::rocprim::block_load_method::block_load_transpose,
820  ::rocprim::block_store_method::block_store_transpose,
821  block_scan_algorithm::using_warp_scan>
822 {};
823 
824 // Based on key_type = short, value_type = short
825 template<class key_type, class value_type>
826 struct default_scan_by_key_config<
827  static_cast<unsigned int>(target_arch::gfx900),
828  key_type,
829  value_type,
830  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
831  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
832  && (sizeof(value_type) > 1))>>
833  : scan_by_key_config_v2<256,
834  22,
835  ::rocprim::block_load_method::block_load_transpose,
836  ::rocprim::block_store_method::block_store_transpose,
837  block_scan_algorithm::reduce_then_scan>
838 {};
839 
840 // Based on key_type = short, value_type = int8_t
841 template<class key_type, class value_type>
842 struct default_scan_by_key_config<
843  static_cast<unsigned int>(target_arch::gfx900),
844  key_type,
845  value_type,
846  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
847  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1))>>
848  : scan_by_key_config_v2<256,
849  20,
850  ::rocprim::block_load_method::block_load_transpose,
851  ::rocprim::block_store_method::block_store_transpose,
852  block_scan_algorithm::using_warp_scan>
853 {};
854 
855 // Based on key_type = int8_t, value_type = int64_t
856 template<class key_type, class value_type>
857 struct default_scan_by_key_config<
858  static_cast<unsigned int>(target_arch::gfx900),
859  key_type,
860  value_type,
861  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
862  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
863  : scan_by_key_config_v2<64,
864  18,
865  ::rocprim::block_load_method::block_load_transpose,
866  ::rocprim::block_store_method::block_store_transpose,
867  block_scan_algorithm::using_warp_scan>
868 {};
869 
870 // Based on key_type = int8_t, value_type = int
871 template<class key_type, class value_type>
872 struct default_scan_by_key_config<
873  static_cast<unsigned int>(target_arch::gfx900),
874  key_type,
875  value_type,
876  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
877  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
878  : scan_by_key_config_v2<256,
879  17,
880  ::rocprim::block_load_method::block_load_transpose,
881  ::rocprim::block_store_method::block_store_transpose,
882  block_scan_algorithm::using_warp_scan>
883 {};
884 
885 // Based on key_type = int8_t, value_type = short
886 template<class key_type, class value_type>
887 struct default_scan_by_key_config<
888  static_cast<unsigned int>(target_arch::gfx900),
889  key_type,
890  value_type,
891  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
892  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
893  : scan_by_key_config_v2<64,
894  24,
895  ::rocprim::block_load_method::block_load_transpose,
896  ::rocprim::block_store_method::block_store_transpose,
897  block_scan_algorithm::reduce_then_scan>
898 {};
899 
900 // Based on key_type = int8_t, value_type = int8_t
901 template<class key_type, class value_type>
902 struct default_scan_by_key_config<
903  static_cast<unsigned int>(target_arch::gfx900),
904  key_type,
905  value_type,
906  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
907  && (sizeof(value_type) <= 1))>>
908  : scan_by_key_config_v2<64,
909  24,
910  ::rocprim::block_load_method::block_load_transpose,
911  ::rocprim::block_store_method::block_store_transpose,
912  block_scan_algorithm::using_warp_scan>
913 {};
914 
915 // Based on key_type = double, value_type = int64_t
916 template<class key_type, class value_type>
917 struct default_scan_by_key_config<
918  static_cast<unsigned int>(target_arch::gfx906),
919  key_type,
920  value_type,
921  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
922  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
923  && (sizeof(value_type) > 4))>>
924  : scan_by_key_config_v2<256,
925  14,
926  ::rocprim::block_load_method::block_load_transpose,
927  ::rocprim::block_store_method::block_store_transpose,
928  block_scan_algorithm::using_warp_scan>
929 {};
930 
931 // Based on key_type = double, value_type = int
932 template<class key_type, class value_type>
933 struct default_scan_by_key_config<
934  static_cast<unsigned int>(target_arch::gfx906),
935  key_type,
936  value_type,
937  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
938  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
939  && (sizeof(value_type) > 2))>>
940  : scan_by_key_config_v2<256,
941  14,
942  ::rocprim::block_load_method::block_load_transpose,
943  ::rocprim::block_store_method::block_store_transpose,
944  block_scan_algorithm::using_warp_scan>
945 {};
946 
947 // Based on key_type = double, value_type = short
948 template<class key_type, class value_type>
949 struct default_scan_by_key_config<
950  static_cast<unsigned int>(target_arch::gfx906),
951  key_type,
952  value_type,
953  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
954  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
955  && (sizeof(value_type) > 1))>>
956  : scan_by_key_config_v2<64,
957  23,
958  ::rocprim::block_load_method::block_load_transpose,
959  ::rocprim::block_store_method::block_store_transpose,
960  block_scan_algorithm::using_warp_scan>
961 {};
962 
963 // Based on key_type = double, value_type = int8_t
964 template<class key_type, class value_type>
965 struct default_scan_by_key_config<
966  static_cast<unsigned int>(target_arch::gfx906),
967  key_type,
968  value_type,
969  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
970  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
971  : scan_by_key_config_v2<64,
972  14,
973  ::rocprim::block_load_method::block_load_transpose,
974  ::rocprim::block_store_method::block_store_transpose,
975  block_scan_algorithm::using_warp_scan>
976 {};
977 
978 // Based on key_type = float, value_type = int64_t
979 template<class key_type, class value_type>
980 struct default_scan_by_key_config<
981  static_cast<unsigned int>(target_arch::gfx906),
982  key_type,
983  value_type,
984  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
985  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
986  && (sizeof(value_type) > 4))>>
987  : scan_by_key_config_v2<64,
988  15,
989  ::rocprim::block_load_method::block_load_transpose,
990  ::rocprim::block_store_method::block_store_transpose,
991  block_scan_algorithm::reduce_then_scan>
992 {};
993 
994 // Based on key_type = float, value_type = int
995 template<class key_type, class value_type>
996 struct default_scan_by_key_config<
997  static_cast<unsigned int>(target_arch::gfx906),
998  key_type,
999  value_type,
1000  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1001  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1002  && (sizeof(value_type) > 2))>>
1003  : scan_by_key_config_v2<256,
1004  15,
1005  ::rocprim::block_load_method::block_load_transpose,
1006  ::rocprim::block_store_method::block_store_transpose,
1007  block_scan_algorithm::reduce_then_scan>
1008 {};
1009 
1010 // Based on key_type = float, value_type = short
1011 template<class key_type, class value_type>
1012 struct default_scan_by_key_config<
1013  static_cast<unsigned int>(target_arch::gfx906),
1014  key_type,
1015  value_type,
1016  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1017  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1018  && (sizeof(value_type) > 1))>>
1019  : scan_by_key_config_v2<256,
1020  14,
1021  ::rocprim::block_load_method::block_load_transpose,
1022  ::rocprim::block_store_method::block_store_transpose,
1023  block_scan_algorithm::using_warp_scan>
1024 {};
1025 
1026 // Based on key_type = float, value_type = int8_t
1027 template<class key_type, class value_type>
1028 struct default_scan_by_key_config<
1029  static_cast<unsigned int>(target_arch::gfx906),
1030  key_type,
1031  value_type,
1032  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1033  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
1034  : scan_by_key_config_v2<64,
1035  20,
1036  ::rocprim::block_load_method::block_load_transpose,
1037  ::rocprim::block_store_method::block_store_transpose,
1038  block_scan_algorithm::reduce_then_scan>
1039 {};
1040 
1041 // Based on key_type = rocprim::half, value_type = int64_t
1042 template<class key_type, class value_type>
1043 struct default_scan_by_key_config<
1044  static_cast<unsigned int>(target_arch::gfx906),
1045  key_type,
1046  value_type,
1047  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1048  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1049  : scan_by_key_config_v2<64,
1050  18,
1051  ::rocprim::block_load_method::block_load_transpose,
1052  ::rocprim::block_store_method::block_store_transpose,
1053  block_scan_algorithm::reduce_then_scan>
1054 {};
1055 
1056 // Based on key_type = rocprim::half, value_type = int
1057 template<class key_type, class value_type>
1058 struct default_scan_by_key_config<
1059  static_cast<unsigned int>(target_arch::gfx906),
1060  key_type,
1061  value_type,
1062  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1063  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1064  : scan_by_key_config_v2<256,
1065  18,
1066  ::rocprim::block_load_method::block_load_transpose,
1067  ::rocprim::block_store_method::block_store_transpose,
1068  block_scan_algorithm::using_warp_scan>
1069 {};
1070 
1071 // Based on key_type = rocprim::half, value_type = short
1072 template<class key_type, class value_type>
1073 struct default_scan_by_key_config<
1074  static_cast<unsigned int>(target_arch::gfx906),
1075  key_type,
1076  value_type,
1077  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1078  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1079  : scan_by_key_config_v2<64,
1080  22,
1081  ::rocprim::block_load_method::block_load_transpose,
1082  ::rocprim::block_store_method::block_store_transpose,
1083  block_scan_algorithm::reduce_then_scan>
1084 {};
1085 
1086 // Based on key_type = rocprim::half, value_type = int8_t
1087 template<class key_type, class value_type>
1088 struct default_scan_by_key_config<
1089  static_cast<unsigned int>(target_arch::gfx906),
1090  key_type,
1091  value_type,
1092  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1093  && (sizeof(value_type) <= 1))>>
1094  : scan_by_key_config_v2<64,
1095  24,
1096  ::rocprim::block_load_method::block_load_transpose,
1097  ::rocprim::block_store_method::block_store_transpose,
1098  block_scan_algorithm::using_warp_scan>
1099 {};
1100 
1101 // Based on key_type = int64_t, value_type = int64_t
1102 template<class key_type, class value_type>
1103 struct default_scan_by_key_config<
1104  static_cast<unsigned int>(target_arch::gfx906),
1105  key_type,
1106  value_type,
1107  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1108  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1109  && (sizeof(value_type) > 4))>>
1110  : scan_by_key_config_v2<256,
1111  14,
1112  ::rocprim::block_load_method::block_load_transpose,
1113  ::rocprim::block_store_method::block_store_transpose,
1114  block_scan_algorithm::using_warp_scan>
1115 {};
1116 
1117 // Based on key_type = int64_t, value_type = int
1118 template<class key_type, class value_type>
1119 struct default_scan_by_key_config<
1120  static_cast<unsigned int>(target_arch::gfx906),
1121  key_type,
1122  value_type,
1123  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1124  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1125  && (sizeof(value_type) > 2))>>
1126  : scan_by_key_config_v2<256,
1127  14,
1128  ::rocprim::block_load_method::block_load_transpose,
1129  ::rocprim::block_store_method::block_store_transpose,
1130  block_scan_algorithm::using_warp_scan>
1131 {};
1132 
1133 // Based on key_type = int64_t, value_type = short
1134 template<class key_type, class value_type>
1135 struct default_scan_by_key_config<
1136  static_cast<unsigned int>(target_arch::gfx906),
1137  key_type,
1138  value_type,
1139  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1140  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
1141  && (sizeof(value_type) > 1))>>
1142  : scan_by_key_config_v2<64,
1143  23,
1144  ::rocprim::block_load_method::block_load_transpose,
1145  ::rocprim::block_store_method::block_store_transpose,
1146  block_scan_algorithm::reduce_then_scan>
1147 {};
1148 
1149 // Based on key_type = int64_t, value_type = int8_t
1150 template<class key_type, class value_type>
1151 struct default_scan_by_key_config<
1152  static_cast<unsigned int>(target_arch::gfx906),
1153  key_type,
1154  value_type,
1155  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1156  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
1157  : scan_by_key_config_v2<64,
1158  14,
1159  ::rocprim::block_load_method::block_load_transpose,
1160  ::rocprim::block_store_method::block_store_transpose,
1161  block_scan_algorithm::using_warp_scan>
1162 {};
1163 
1164 // Based on key_type = int, value_type = int64_t
1165 template<class key_type, class value_type>
1166 struct default_scan_by_key_config<
1167  static_cast<unsigned int>(target_arch::gfx906),
1168  key_type,
1169  value_type,
1170  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1171  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
1172  && (sizeof(value_type) > 4))>>
1173  : scan_by_key_config_v2<64,
1174  15,
1175  ::rocprim::block_load_method::block_load_transpose,
1176  ::rocprim::block_store_method::block_store_transpose,
1177  block_scan_algorithm::reduce_then_scan>
1178 {};
1179 
1180 // Based on key_type = int, value_type = int
1181 template<class key_type, class value_type>
1182 struct default_scan_by_key_config<
1183  static_cast<unsigned int>(target_arch::gfx906),
1184  key_type,
1185  value_type,
1186  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1187  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1188  && (sizeof(value_type) > 2))>>
1189  : scan_by_key_config_v2<256,
1190  15,
1191  ::rocprim::block_load_method::block_load_transpose,
1192  ::rocprim::block_store_method::block_store_transpose,
1193  block_scan_algorithm::reduce_then_scan>
1194 {};
1195 
1196 // Based on key_type = int, value_type = short
1197 template<class key_type, class value_type>
1198 struct default_scan_by_key_config<
1199  static_cast<unsigned int>(target_arch::gfx906),
1200  key_type,
1201  value_type,
1202  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1203  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1204  && (sizeof(value_type) > 1))>>
1205  : scan_by_key_config_v2<256,
1206  14,
1207  ::rocprim::block_load_method::block_load_transpose,
1208  ::rocprim::block_store_method::block_store_transpose,
1209  block_scan_algorithm::using_warp_scan>
1210 {};
1211 
1212 // Based on key_type = int, value_type = int8_t
1213 template<class key_type, class value_type>
1214 struct default_scan_by_key_config<
1215  static_cast<unsigned int>(target_arch::gfx906),
1216  key_type,
1217  value_type,
1218  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1219  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
1220  : scan_by_key_config_v2<64,
1221  20,
1222  ::rocprim::block_load_method::block_load_transpose,
1223  ::rocprim::block_store_method::block_store_transpose,
1224  block_scan_algorithm::using_warp_scan>
1225 {};
1226 
1227 // Based on key_type = short, value_type = int64_t
1228 template<class key_type, class value_type>
1229 struct default_scan_by_key_config<
1230  static_cast<unsigned int>(target_arch::gfx906),
1231  key_type,
1232  value_type,
1233  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1234  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
1235  && (sizeof(value_type) > 4))>>
1236  : scan_by_key_config_v2<64,
1237  18,
1238  ::rocprim::block_load_method::block_load_transpose,
1239  ::rocprim::block_store_method::block_store_transpose,
1240  block_scan_algorithm::using_warp_scan>
1241 {};
1242 
1243 // Based on key_type = short, value_type = int
1244 template<class key_type, class value_type>
1245 struct default_scan_by_key_config<
1246  static_cast<unsigned int>(target_arch::gfx906),
1247  key_type,
1248  value_type,
1249  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1250  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
1251  && (sizeof(value_type) > 2))>>
1252  : scan_by_key_config_v2<256,
1253  18,
1254  ::rocprim::block_load_method::block_load_transpose,
1255  ::rocprim::block_store_method::block_store_transpose,
1256  block_scan_algorithm::using_warp_scan>
1257 {};
1258 
1259 // Based on key_type = short, value_type = short
1260 template<class key_type, class value_type>
1261 struct default_scan_by_key_config<
1262  static_cast<unsigned int>(target_arch::gfx906),
1263  key_type,
1264  value_type,
1265  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1266  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
1267  && (sizeof(value_type) > 1))>>
1268  : scan_by_key_config_v2<64,
1269  22,
1270  ::rocprim::block_load_method::block_load_transpose,
1271  ::rocprim::block_store_method::block_store_transpose,
1272  block_scan_algorithm::using_warp_scan>
1273 {};
1274 
1275 // Based on key_type = short, value_type = int8_t
1276 template<class key_type, class value_type>
1277 struct default_scan_by_key_config<
1278  static_cast<unsigned int>(target_arch::gfx906),
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(key_type) > 1) && (sizeof(value_type) <= 1))>>
1283  : scan_by_key_config_v2<64,
1284  24,
1285  ::rocprim::block_load_method::block_load_transpose,
1286  ::rocprim::block_store_method::block_store_transpose,
1287  block_scan_algorithm::using_warp_scan>
1288 {};
1289 
1290 // Based on key_type = int8_t, value_type = int64_t
1291 template<class key_type, class value_type>
1292 struct default_scan_by_key_config<
1293  static_cast<unsigned int>(target_arch::gfx906),
1294  key_type,
1295  value_type,
1296  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1297  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1298  : scan_by_key_config_v2<64,
1299  18,
1300  ::rocprim::block_load_method::block_load_transpose,
1301  ::rocprim::block_store_method::block_store_transpose,
1302  block_scan_algorithm::reduce_then_scan>
1303 {};
1304 
1305 // Based on key_type = int8_t, value_type = int
1306 template<class key_type, class value_type>
1307 struct default_scan_by_key_config<
1308  static_cast<unsigned int>(target_arch::gfx906),
1309  key_type,
1310  value_type,
1311  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1312  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1313  : scan_by_key_config_v2<256,
1314  13,
1315  ::rocprim::block_load_method::block_load_transpose,
1316  ::rocprim::block_store_method::block_store_transpose,
1317  block_scan_algorithm::using_warp_scan>
1318 {};
1319 
1320 // Based on key_type = int8_t, value_type = short
1321 template<class key_type, class value_type>
1322 struct default_scan_by_key_config<
1323  static_cast<unsigned int>(target_arch::gfx906),
1324  key_type,
1325  value_type,
1326  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1327  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1328  : scan_by_key_config_v2<64,
1329  24,
1330  ::rocprim::block_load_method::block_load_transpose,
1331  ::rocprim::block_store_method::block_store_transpose,
1332  block_scan_algorithm::using_warp_scan>
1333 {};
1334 
1335 // Based on key_type = int8_t, value_type = int8_t
1336 template<class key_type, class value_type>
1337 struct default_scan_by_key_config<
1338  static_cast<unsigned int>(target_arch::gfx906),
1339  key_type,
1340  value_type,
1341  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1342  && (sizeof(value_type) <= 1))>>
1343  : scan_by_key_config_v2<64,
1344  24,
1345  ::rocprim::block_load_method::block_load_transpose,
1346  ::rocprim::block_store_method::block_store_transpose,
1347  block_scan_algorithm::reduce_then_scan>
1348 {};
1349 
1350 // Based on key_type = double, value_type = int64_t
1351 template<class key_type, class value_type>
1352 struct default_scan_by_key_config<
1353  static_cast<unsigned int>(target_arch::gfx1030),
1354  key_type,
1355  value_type,
1356  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1357  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1358  && (sizeof(value_type) > 4))>>
1359  : scan_by_key_config_v2<128,
1360  23,
1361  ::rocprim::block_load_method::block_load_transpose,
1362  ::rocprim::block_store_method::block_store_transpose,
1363  block_scan_algorithm::using_warp_scan>
1364 {};
1365 
1366 // Based on key_type = double, value_type = int
1367 template<class key_type, class value_type>
1368 struct default_scan_by_key_config<
1369  static_cast<unsigned int>(target_arch::gfx1030),
1370  key_type,
1371  value_type,
1372  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1373  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1374  && (sizeof(value_type) > 2))>>
1375  : scan_by_key_config_v2<128,
1376  23,
1377  ::rocprim::block_load_method::block_load_transpose,
1378  ::rocprim::block_store_method::block_store_transpose,
1379  block_scan_algorithm::reduce_then_scan>
1380 {};
1381 
1382 // Based on key_type = double, value_type = short
1383 template<class key_type, class value_type>
1384 struct default_scan_by_key_config<
1385  static_cast<unsigned int>(target_arch::gfx1030),
1386  key_type,
1387  value_type,
1388  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1389  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
1390  && (sizeof(value_type) > 1))>>
1391  : scan_by_key_config_v2<64,
1392  22,
1393  ::rocprim::block_load_method::block_load_transpose,
1394  ::rocprim::block_store_method::block_store_transpose,
1395  block_scan_algorithm::using_warp_scan>
1396 {};
1397 
1398 // Based on key_type = double, value_type = int8_t
1399 template<class key_type, class value_type>
1400 struct default_scan_by_key_config<
1401  static_cast<unsigned int>(target_arch::gfx1030),
1402  key_type,
1403  value_type,
1404  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1405  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
1406  : scan_by_key_config_v2<64,
1407  17,
1408  ::rocprim::block_load_method::block_load_transpose,
1409  ::rocprim::block_store_method::block_store_transpose,
1410  block_scan_algorithm::using_warp_scan>
1411 {};
1412 
1413 // Based on key_type = float, value_type = int64_t
1414 template<class key_type, class value_type>
1415 struct default_scan_by_key_config<
1416  static_cast<unsigned int>(target_arch::gfx1030),
1417  key_type,
1418  value_type,
1419  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1420  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
1421  && (sizeof(value_type) > 4))>>
1422  : scan_by_key_config_v2<256,
1423  9,
1424  ::rocprim::block_load_method::block_load_transpose,
1425  ::rocprim::block_store_method::block_store_transpose,
1426  block_scan_algorithm::using_warp_scan>
1427 {};
1428 
1429 // Based on key_type = float, value_type = int
1430 template<class key_type, class value_type>
1431 struct default_scan_by_key_config<
1432  static_cast<unsigned int>(target_arch::gfx1030),
1433  key_type,
1434  value_type,
1435  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1436  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1437  && (sizeof(value_type) > 2))>>
1438  : scan_by_key_config_v2<256,
1439  15,
1440  ::rocprim::block_load_method::block_load_transpose,
1441  ::rocprim::block_store_method::block_store_transpose,
1442  block_scan_algorithm::using_warp_scan>
1443 {};
1444 
1445 // Based on key_type = float, value_type = short
1446 template<class key_type, class value_type>
1447 struct default_scan_by_key_config<
1448  static_cast<unsigned int>(target_arch::gfx1030),
1449  key_type,
1450  value_type,
1451  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1452  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1453  && (sizeof(value_type) > 1))>>
1454  : scan_by_key_config_v2<64,
1455  22,
1456  ::rocprim::block_load_method::block_load_transpose,
1457  ::rocprim::block_store_method::block_store_transpose,
1458  block_scan_algorithm::reduce_then_scan>
1459 {};
1460 
1461 // Based on key_type = float, value_type = int8_t
1462 template<class key_type, class value_type>
1463 struct default_scan_by_key_config<
1464  static_cast<unsigned int>(target_arch::gfx1030),
1465  key_type,
1466  value_type,
1467  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1468  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
1469  : scan_by_key_config_v2<128,
1470  7,
1471  ::rocprim::block_load_method::block_load_transpose,
1472  ::rocprim::block_store_method::block_store_transpose,
1473  block_scan_algorithm::reduce_then_scan>
1474 {};
1475 
1476 // Based on key_type = rocprim::half, value_type = int64_t
1477 template<class key_type, class value_type>
1478 struct default_scan_by_key_config<
1479  static_cast<unsigned int>(target_arch::gfx1030),
1480  key_type,
1481  value_type,
1482  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1483  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1484  : scan_by_key_config_v2<64,
1485  23,
1486  ::rocprim::block_load_method::block_load_transpose,
1487  ::rocprim::block_store_method::block_store_transpose,
1488  block_scan_algorithm::using_warp_scan>
1489 {};
1490 
1491 // Based on key_type = rocprim::half, value_type = int
1492 template<class key_type, class value_type>
1493 struct default_scan_by_key_config<
1494  static_cast<unsigned int>(target_arch::gfx1030),
1495  key_type,
1496  value_type,
1497  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1498  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1499  : scan_by_key_config_v2<128,
1500  22,
1501  ::rocprim::block_load_method::block_load_transpose,
1502  ::rocprim::block_store_method::block_store_transpose,
1503  block_scan_algorithm::using_warp_scan>
1504 {};
1505 
1506 // Based on key_type = rocprim::half, value_type = short
1507 template<class key_type, class value_type>
1508 struct default_scan_by_key_config<
1509  static_cast<unsigned int>(target_arch::gfx1030),
1510  key_type,
1511  value_type,
1512  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1513  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1514  : scan_by_key_config_v2<128,
1515  22,
1516  ::rocprim::block_load_method::block_load_transpose,
1517  ::rocprim::block_store_method::block_store_transpose,
1518  block_scan_algorithm::using_warp_scan>
1519 {};
1520 
1521 // Based on key_type = rocprim::half, value_type = int8_t
1522 template<class key_type, class value_type>
1523 struct default_scan_by_key_config<
1524  static_cast<unsigned int>(target_arch::gfx1030),
1525  key_type,
1526  value_type,
1527  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1528  && (sizeof(value_type) <= 1))>>
1529  : scan_by_key_config_v2<256,
1530  20,
1531  ::rocprim::block_load_method::block_load_transpose,
1532  ::rocprim::block_store_method::block_store_transpose,
1533  block_scan_algorithm::using_warp_scan>
1534 {};
1535 
1536 // Based on key_type = int64_t, value_type = int64_t
1537 template<class key_type, class value_type>
1538 struct default_scan_by_key_config<
1539  static_cast<unsigned int>(target_arch::gfx1030),
1540  key_type,
1541  value_type,
1542  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1543  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1544  && (sizeof(value_type) > 4))>>
1545  : scan_by_key_config_v2<256,
1546  13,
1547  ::rocprim::block_load_method::block_load_transpose,
1548  ::rocprim::block_store_method::block_store_transpose,
1549  block_scan_algorithm::using_warp_scan>
1550 {};
1551 
1552 // Based on key_type = int64_t, value_type = int
1553 template<class key_type, class value_type>
1554 struct default_scan_by_key_config<
1555  static_cast<unsigned int>(target_arch::gfx1030),
1556  key_type,
1557  value_type,
1558  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1559  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1560  && (sizeof(value_type) > 2))>>
1561  : scan_by_key_config_v2<128,
1562  19,
1563  ::rocprim::block_load_method::block_load_transpose,
1564  ::rocprim::block_store_method::block_store_transpose,
1565  block_scan_algorithm::using_warp_scan>
1566 {};
1567 
1568 // Based on key_type = int64_t, value_type = short
1569 template<class key_type, class value_type>
1570 struct default_scan_by_key_config<
1571  static_cast<unsigned int>(target_arch::gfx1030),
1572  key_type,
1573  value_type,
1574  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1575  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
1576  && (sizeof(value_type) > 1))>>
1577  : scan_by_key_config_v2<64,
1578  22,
1579  ::rocprim::block_load_method::block_load_transpose,
1580  ::rocprim::block_store_method::block_store_transpose,
1581  block_scan_algorithm::using_warp_scan>
1582 {};
1583 
1584 // Based on key_type = int64_t, value_type = int8_t
1585 template<class key_type, class value_type>
1586 struct default_scan_by_key_config<
1587  static_cast<unsigned int>(target_arch::gfx1030),
1588  key_type,
1589  value_type,
1590  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1591  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
1592  : scan_by_key_config_v2<64,
1593  18,
1594  ::rocprim::block_load_method::block_load_transpose,
1595  ::rocprim::block_store_method::block_store_transpose,
1596  block_scan_algorithm::reduce_then_scan>
1597 {};
1598 
1599 // Based on key_type = int, value_type = int64_t
1600 template<class key_type, class value_type>
1601 struct default_scan_by_key_config<
1602  static_cast<unsigned int>(target_arch::gfx1030),
1603  key_type,
1604  value_type,
1605  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1606  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
1607  && (sizeof(value_type) > 4))>>
1608  : scan_by_key_config_v2<256,
1609  9,
1610  ::rocprim::block_load_method::block_load_transpose,
1611  ::rocprim::block_store_method::block_store_transpose,
1612  block_scan_algorithm::using_warp_scan>
1613 {};
1614 
1615 // Based on key_type = int, value_type = int
1616 template<class key_type, class value_type>
1617 struct default_scan_by_key_config<
1618  static_cast<unsigned int>(target_arch::gfx1030),
1619  key_type,
1620  value_type,
1621  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1622  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1623  && (sizeof(value_type) > 2))>>
1624  : scan_by_key_config_v2<128,
1625  22,
1626  ::rocprim::block_load_method::block_load_transpose,
1627  ::rocprim::block_store_method::block_store_transpose,
1628  block_scan_algorithm::reduce_then_scan>
1629 {};
1630 
1631 // Based on key_type = int, value_type = short
1632 template<class key_type, class value_type>
1633 struct default_scan_by_key_config<
1634  static_cast<unsigned int>(target_arch::gfx1030),
1635  key_type,
1636  value_type,
1637  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1638  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1639  && (sizeof(value_type) > 1))>>
1640  : scan_by_key_config_v2<64,
1641  22,
1642  ::rocprim::block_load_method::block_load_transpose,
1643  ::rocprim::block_store_method::block_store_transpose,
1644  block_scan_algorithm::reduce_then_scan>
1645 {};
1646 
1647 // Based on key_type = int, value_type = int8_t
1648 template<class key_type, class value_type>
1649 struct default_scan_by_key_config<
1650  static_cast<unsigned int>(target_arch::gfx1030),
1651  key_type,
1652  value_type,
1653  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1654  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
1655  : scan_by_key_config_v2<128,
1656  7,
1657  ::rocprim::block_load_method::block_load_transpose,
1658  ::rocprim::block_store_method::block_store_transpose,
1659  block_scan_algorithm::reduce_then_scan>
1660 {};
1661 
1662 // Based on key_type = short, value_type = int64_t
1663 template<class key_type, class value_type>
1664 struct default_scan_by_key_config<
1665  static_cast<unsigned int>(target_arch::gfx1030),
1666  key_type,
1667  value_type,
1668  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1669  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
1670  && (sizeof(value_type) > 4))>>
1671  : scan_by_key_config_v2<256,
1672  9,
1673  ::rocprim::block_load_method::block_load_transpose,
1674  ::rocprim::block_store_method::block_store_transpose,
1675  block_scan_algorithm::using_warp_scan>
1676 {};
1677 
1678 // Based on key_type = short, value_type = int
1679 template<class key_type, class value_type>
1680 struct default_scan_by_key_config<
1681  static_cast<unsigned int>(target_arch::gfx1030),
1682  key_type,
1683  value_type,
1684  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1685  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
1686  && (sizeof(value_type) > 2))>>
1687  : scan_by_key_config_v2<128,
1688  22,
1689  ::rocprim::block_load_method::block_load_transpose,
1690  ::rocprim::block_store_method::block_store_transpose,
1691  block_scan_algorithm::using_warp_scan>
1692 {};
1693 
1694 // Based on key_type = short, value_type = short
1695 template<class key_type, class value_type>
1696 struct default_scan_by_key_config<
1697  static_cast<unsigned int>(target_arch::gfx1030),
1698  key_type,
1699  value_type,
1700  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1701  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
1702  && (sizeof(value_type) > 1))>>
1703  : scan_by_key_config_v2<128,
1704  22,
1705  ::rocprim::block_load_method::block_load_transpose,
1706  ::rocprim::block_store_method::block_store_transpose,
1707  block_scan_algorithm::using_warp_scan>
1708 {};
1709 
1710 // Based on key_type = short, value_type = int8_t
1711 template<class key_type, class value_type>
1712 struct default_scan_by_key_config<
1713  static_cast<unsigned int>(target_arch::gfx1030),
1714  key_type,
1715  value_type,
1716  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
1717  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1))>>
1718  : scan_by_key_config_v2<256,
1719  20,
1720  ::rocprim::block_load_method::block_load_transpose,
1721  ::rocprim::block_store_method::block_store_transpose,
1722  block_scan_algorithm::using_warp_scan>
1723 {};
1724 
1725 // Based on key_type = int8_t, value_type = int64_t
1726 template<class key_type, class value_type>
1727 struct default_scan_by_key_config<
1728  static_cast<unsigned int>(target_arch::gfx1030),
1729  key_type,
1730  value_type,
1731  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1732  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1733  : scan_by_key_config_v2<128,
1734  17,
1735  ::rocprim::block_load_method::block_load_transpose,
1736  ::rocprim::block_store_method::block_store_transpose,
1737  block_scan_algorithm::reduce_then_scan>
1738 {};
1739 
1740 // Based on key_type = int8_t, value_type = int
1741 template<class key_type, class value_type>
1742 struct default_scan_by_key_config<
1743  static_cast<unsigned int>(target_arch::gfx1030),
1744  key_type,
1745  value_type,
1746  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1747  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1748  : scan_by_key_config_v2<256,
1749  12,
1750  ::rocprim::block_load_method::block_load_transpose,
1751  ::rocprim::block_store_method::block_store_transpose,
1752  block_scan_algorithm::using_warp_scan>
1753 {};
1754 
1755 // Based on key_type = int8_t, value_type = short
1756 template<class key_type, class value_type>
1757 struct default_scan_by_key_config<
1758  static_cast<unsigned int>(target_arch::gfx1030),
1759  key_type,
1760  value_type,
1761  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1762  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1763  : scan_by_key_config_v2<128,
1764  24,
1765  ::rocprim::block_load_method::block_load_transpose,
1766  ::rocprim::block_store_method::block_store_transpose,
1767  block_scan_algorithm::using_warp_scan>
1768 {};
1769 
1770 // Based on key_type = int8_t, value_type = int8_t
1771 template<class key_type, class value_type>
1772 struct default_scan_by_key_config<
1773  static_cast<unsigned int>(target_arch::gfx1030),
1774  key_type,
1775  value_type,
1776  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
1777  && (sizeof(value_type) <= 1))>>
1778  : scan_by_key_config_v2<256,
1779  20,
1780  ::rocprim::block_load_method::block_load_transpose,
1781  ::rocprim::block_store_method::block_store_transpose,
1782  block_scan_algorithm::using_warp_scan>
1783 {};
1784 
1785 // Based on key_type = double, value_type = int64_t
1786 template<class key_type, class value_type>
1787 struct default_scan_by_key_config<
1788  static_cast<unsigned int>(target_arch::unknown),
1789  key_type,
1790  value_type,
1791  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1792  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1793  && (sizeof(value_type) > 4))>>
1794  : scan_by_key_config_v2<256,
1795  14,
1796  ::rocprim::block_load_method::block_load_transpose,
1797  ::rocprim::block_store_method::block_store_transpose,
1798  block_scan_algorithm::using_warp_scan>
1799 {};
1800 
1801 // Based on key_type = double, value_type = int
1802 template<class key_type, class value_type>
1803 struct default_scan_by_key_config<
1804  static_cast<unsigned int>(target_arch::unknown),
1805  key_type,
1806  value_type,
1807  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1808  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1809  && (sizeof(value_type) > 2))>>
1810  : scan_by_key_config_v2<256,
1811  14,
1812  ::rocprim::block_load_method::block_load_transpose,
1813  ::rocprim::block_store_method::block_store_transpose,
1814  block_scan_algorithm::reduce_then_scan>
1815 {};
1816 
1817 // Based on key_type = double, value_type = short
1818 template<class key_type, class value_type>
1819 struct default_scan_by_key_config<
1820  static_cast<unsigned int>(target_arch::unknown),
1821  key_type,
1822  value_type,
1823  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1824  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
1825  && (sizeof(value_type) > 1))>>
1826  : scan_by_key_config_v2<64,
1827  14,
1828  ::rocprim::block_load_method::block_load_transpose,
1829  ::rocprim::block_store_method::block_store_transpose,
1830  block_scan_algorithm::using_warp_scan>
1831 {};
1832 
1833 // Based on key_type = double, value_type = int8_t
1834 template<class key_type, class value_type>
1835 struct default_scan_by_key_config<
1836  static_cast<unsigned int>(target_arch::unknown),
1837  key_type,
1838  value_type,
1839  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1840  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
1841  : scan_by_key_config_v2<64,
1842  14,
1843  ::rocprim::block_load_method::block_load_transpose,
1844  ::rocprim::block_store_method::block_store_transpose,
1845  block_scan_algorithm::reduce_then_scan>
1846 {};
1847 
1848 // Based on key_type = float, value_type = int64_t
1849 template<class key_type, class value_type>
1850 struct default_scan_by_key_config<
1851  static_cast<unsigned int>(target_arch::unknown),
1852  key_type,
1853  value_type,
1854  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1855  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
1856  && (sizeof(value_type) > 4))>>
1857  : scan_by_key_config_v2<128,
1858  10,
1859  ::rocprim::block_load_method::block_load_transpose,
1860  ::rocprim::block_store_method::block_store_transpose,
1861  block_scan_algorithm::using_warp_scan>
1862 {};
1863 
1864 // Based on key_type = float, value_type = int
1865 template<class key_type, class value_type>
1866 struct default_scan_by_key_config<
1867  static_cast<unsigned int>(target_arch::unknown),
1868  key_type,
1869  value_type,
1870  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1871  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
1872  && (sizeof(value_type) > 2))>>
1873  : scan_by_key_config_v2<256,
1874  12,
1875  ::rocprim::block_load_method::block_load_transpose,
1876  ::rocprim::block_store_method::block_store_transpose,
1877  block_scan_algorithm::reduce_then_scan>
1878 {};
1879 
1880 // Based on key_type = float, value_type = short
1881 template<class key_type, class value_type>
1882 struct default_scan_by_key_config<
1883  static_cast<unsigned int>(target_arch::unknown),
1884  key_type,
1885  value_type,
1886  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
1887  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
1888  && (sizeof(value_type) > 1))>>
1889  : scan_by_key_config_v2<256,
1890  14,
1891  ::rocprim::block_load_method::block_load_transpose,
1892  ::rocprim::block_store_method::block_store_transpose,
1893  block_scan_algorithm::reduce_then_scan>
1894 {};
1895 
1896 // Based on key_type = float, value_type = int8_t
1897 template<class key_type, class value_type>
1898 struct default_scan_by_key_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) <= 1))>>
1904  : scan_by_key_config_v2<128,
1905  14,
1906  ::rocprim::block_load_method::block_load_transpose,
1907  ::rocprim::block_store_method::block_store_transpose,
1908  block_scan_algorithm::using_warp_scan>
1909 {};
1910 
1911 // Based on key_type = rocprim::half, value_type = int64_t
1912 template<class key_type, class value_type>
1913 struct default_scan_by_key_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) <= 2)
1918  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
1919  : scan_by_key_config_v2<64,
1920  18,
1921  ::rocprim::block_load_method::block_load_transpose,
1922  ::rocprim::block_store_method::block_store_transpose,
1923  block_scan_algorithm::using_warp_scan>
1924 {};
1925 
1926 // Based on key_type = rocprim::half, value_type = int
1927 template<class key_type, class value_type>
1928 struct default_scan_by_key_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) <= 2)
1933  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
1934  : scan_by_key_config_v2<256,
1935  18,
1936  ::rocprim::block_load_method::block_load_transpose,
1937  ::rocprim::block_store_method::block_store_transpose,
1938  block_scan_algorithm::using_warp_scan>
1939 {};
1940 
1941 // Based on key_type = rocprim::half, value_type = short
1942 template<class key_type, class value_type>
1943 struct default_scan_by_key_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) <= 2)
1948  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
1949  : scan_by_key_config_v2<256,
1950  20,
1951  ::rocprim::block_load_method::block_load_transpose,
1952  ::rocprim::block_store_method::block_store_transpose,
1953  block_scan_algorithm::reduce_then_scan>
1954 {};
1955 
1956 // Based on key_type = rocprim::half, value_type = int8_t
1957 template<class key_type, class value_type>
1958 struct default_scan_by_key_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) <= 2)
1963  && (sizeof(value_type) <= 1))>>
1964  : scan_by_key_config_v2<256,
1965  24,
1966  ::rocprim::block_load_method::block_load_transpose,
1967  ::rocprim::block_store_method::block_store_transpose,
1968  block_scan_algorithm::using_warp_scan>
1969 {};
1970 
1971 // Based on key_type = int64_t, value_type = int64_t
1972 template<class key_type, class value_type>
1973 struct default_scan_by_key_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) <= 8)
1978  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
1979  && (sizeof(value_type) > 4))>>
1980  : scan_by_key_config_v2<256,
1981  14,
1982  ::rocprim::block_load_method::block_load_transpose,
1983  ::rocprim::block_store_method::block_store_transpose,
1984  block_scan_algorithm::reduce_then_scan>
1985 {};
1986 
1987 // Based on key_type = int64_t, value_type = int
1988 template<class key_type, class value_type>
1989 struct default_scan_by_key_config<
1990  static_cast<unsigned int>(target_arch::unknown),
1991  key_type,
1992  value_type,
1993  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
1994  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
1995  && (sizeof(value_type) > 2))>>
1996  : scan_by_key_config_v2<256,
1997  14,
1998  ::rocprim::block_load_method::block_load_transpose,
1999  ::rocprim::block_store_method::block_store_transpose,
2000  block_scan_algorithm::reduce_then_scan>
2001 {};
2002 
2003 // Based on key_type = int64_t, value_type = short
2004 template<class key_type, class value_type>
2005 struct default_scan_by_key_config<
2006  static_cast<unsigned int>(target_arch::unknown),
2007  key_type,
2008  value_type,
2009  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2010  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
2011  && (sizeof(value_type) > 1))>>
2012  : scan_by_key_config_v2<64,
2013  14,
2014  ::rocprim::block_load_method::block_load_transpose,
2015  ::rocprim::block_store_method::block_store_transpose,
2016  block_scan_algorithm::using_warp_scan>
2017 {};
2018 
2019 // Based on key_type = int64_t, value_type = int8_t
2020 template<class key_type, class value_type>
2021 struct default_scan_by_key_config<
2022  static_cast<unsigned int>(target_arch::unknown),
2023  key_type,
2024  value_type,
2025  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2026  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
2027  : scan_by_key_config_v2<256,
2028  14,
2029  ::rocprim::block_load_method::block_load_transpose,
2030  ::rocprim::block_store_method::block_store_transpose,
2031  block_scan_algorithm::using_warp_scan>
2032 {};
2033 
2034 // Based on key_type = int, value_type = int64_t
2035 template<class key_type, class value_type>
2036 struct default_scan_by_key_config<
2037  static_cast<unsigned int>(target_arch::unknown),
2038  key_type,
2039  value_type,
2040  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2041  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
2042  && (sizeof(value_type) > 4))>>
2043  : scan_by_key_config_v2<128,
2044  10,
2045  ::rocprim::block_load_method::block_load_transpose,
2046  ::rocprim::block_store_method::block_store_transpose,
2047  block_scan_algorithm::using_warp_scan>
2048 {};
2049 
2050 // Based on key_type = int, value_type = int
2051 template<class key_type, class value_type>
2052 struct default_scan_by_key_config<
2053  static_cast<unsigned int>(target_arch::unknown),
2054  key_type,
2055  value_type,
2056  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2057  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
2058  && (sizeof(value_type) > 2))>>
2059  : scan_by_key_config_v2<256,
2060  12,
2061  ::rocprim::block_load_method::block_load_transpose,
2062  ::rocprim::block_store_method::block_store_transpose,
2063  block_scan_algorithm::reduce_then_scan>
2064 {};
2065 
2066 // Based on key_type = int, value_type = short
2067 template<class key_type, class value_type>
2068 struct default_scan_by_key_config<
2069  static_cast<unsigned int>(target_arch::unknown),
2070  key_type,
2071  value_type,
2072  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2073  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
2074  && (sizeof(value_type) > 1))>>
2075  : scan_by_key_config_v2<256,
2076  14,
2077  ::rocprim::block_load_method::block_load_transpose,
2078  ::rocprim::block_store_method::block_store_transpose,
2079  block_scan_algorithm::reduce_then_scan>
2080 {};
2081 
2082 // Based on key_type = int, value_type = int8_t
2083 template<class key_type, class value_type>
2084 struct default_scan_by_key_config<
2085  static_cast<unsigned int>(target_arch::unknown),
2086  key_type,
2087  value_type,
2088  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2089  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
2090  : scan_by_key_config_v2<128,
2091  14,
2092  ::rocprim::block_load_method::block_load_transpose,
2093  ::rocprim::block_store_method::block_store_transpose,
2094  block_scan_algorithm::using_warp_scan>
2095 {};
2096 
2097 // Based on key_type = short, value_type = int64_t
2098 template<class key_type, class value_type>
2099 struct default_scan_by_key_config<
2100  static_cast<unsigned int>(target_arch::unknown),
2101  key_type,
2102  value_type,
2103  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2104  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
2105  && (sizeof(value_type) > 4))>>
2106  : scan_by_key_config_v2<64,
2107  18,
2108  ::rocprim::block_load_method::block_load_transpose,
2109  ::rocprim::block_store_method::block_store_transpose,
2110  block_scan_algorithm::reduce_then_scan>
2111 {};
2112 
2113 // Based on key_type = short, value_type = int
2114 template<class key_type, class value_type>
2115 struct default_scan_by_key_config<
2116  static_cast<unsigned int>(target_arch::unknown),
2117  key_type,
2118  value_type,
2119  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2120  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
2121  && (sizeof(value_type) > 2))>>
2122  : scan_by_key_config_v2<256,
2123  18,
2124  ::rocprim::block_load_method::block_load_transpose,
2125  ::rocprim::block_store_method::block_store_transpose,
2126  block_scan_algorithm::using_warp_scan>
2127 {};
2128 
2129 // Based on key_type = short, value_type = short
2130 template<class key_type, class value_type>
2131 struct default_scan_by_key_config<
2132  static_cast<unsigned int>(target_arch::unknown),
2133  key_type,
2134  value_type,
2135  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2136  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
2137  && (sizeof(value_type) > 1))>>
2138  : scan_by_key_config_v2<256,
2139  20,
2140  ::rocprim::block_load_method::block_load_transpose,
2141  ::rocprim::block_store_method::block_store_transpose,
2142  block_scan_algorithm::reduce_then_scan>
2143 {};
2144 
2145 // Based on key_type = short, value_type = int8_t
2146 template<class key_type, class value_type>
2147 struct default_scan_by_key_config<
2148  static_cast<unsigned int>(target_arch::unknown),
2149  key_type,
2150  value_type,
2151  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2152  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1))>>
2153  : scan_by_key_config_v2<256,
2154  24,
2155  ::rocprim::block_load_method::block_load_transpose,
2156  ::rocprim::block_store_method::block_store_transpose,
2157  block_scan_algorithm::using_warp_scan>
2158 {};
2159 
2160 // Based on key_type = int8_t, value_type = int64_t
2161 template<class key_type, class value_type>
2162 struct default_scan_by_key_config<
2163  static_cast<unsigned int>(target_arch::unknown),
2164  key_type,
2165  value_type,
2166  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2167  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
2168  : scan_by_key_config_v2<64,
2169  18,
2170  ::rocprim::block_load_method::block_load_transpose,
2171  ::rocprim::block_store_method::block_store_transpose,
2172  block_scan_algorithm::using_warp_scan>
2173 {};
2174 
2175 // Based on key_type = int8_t, value_type = int
2176 template<class key_type, class value_type>
2177 struct default_scan_by_key_config<
2178  static_cast<unsigned int>(target_arch::unknown),
2179  key_type,
2180  value_type,
2181  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2182  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
2183  : scan_by_key_config_v2<128,
2184  18,
2185  ::rocprim::block_load_method::block_load_transpose,
2186  ::rocprim::block_store_method::block_store_transpose,
2187  block_scan_algorithm::using_warp_scan>
2188 {};
2189 
2190 // Based on key_type = int8_t, value_type = short
2191 template<class key_type, class value_type>
2192 struct default_scan_by_key_config<
2193  static_cast<unsigned int>(target_arch::unknown),
2194  key_type,
2195  value_type,
2196  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2197  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
2198  : scan_by_key_config_v2<256,
2199  20,
2200  ::rocprim::block_load_method::block_load_transpose,
2201  ::rocprim::block_store_method::block_store_transpose,
2202  block_scan_algorithm::reduce_then_scan>
2203 {};
2204 
2205 // Based on key_type = int8_t, value_type = int8_t
2206 template<class key_type, class value_type>
2207 struct default_scan_by_key_config<
2208  static_cast<unsigned int>(target_arch::unknown),
2209  key_type,
2210  value_type,
2211  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2212  && (sizeof(value_type) <= 1))>>
2213  : scan_by_key_config_v2<256,
2214  24,
2215  ::rocprim::block_load_method::block_load_transpose,
2216  ::rocprim::block_store_method::block_store_transpose,
2217  block_scan_algorithm::reduce_then_scan>
2218 {};
2219 
2220 // Based on key_type = double, value_type = int64_t
2221 template<class key_type, class value_type>
2222 struct default_scan_by_key_config<
2223  static_cast<unsigned int>(target_arch::gfx90a),
2224  key_type,
2225  value_type,
2226  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2227  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
2228  && (sizeof(value_type) > 4))>>
2229  : scan_by_key_config_v2<256,
2230  14,
2231  ::rocprim::block_load_method::block_load_transpose,
2232  ::rocprim::block_store_method::block_store_transpose,
2233  block_scan_algorithm::using_warp_scan>
2234 {};
2235 
2236 // Based on key_type = double, value_type = int
2237 template<class key_type, class value_type>
2238 struct default_scan_by_key_config<
2239  static_cast<unsigned int>(target_arch::gfx90a),
2240  key_type,
2241  value_type,
2242  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2243  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
2244  && (sizeof(value_type) > 2))>>
2245  : scan_by_key_config_v2<256,
2246  14,
2247  ::rocprim::block_load_method::block_load_transpose,
2248  ::rocprim::block_store_method::block_store_transpose,
2249  block_scan_algorithm::reduce_then_scan>
2250 {};
2251 
2252 // Based on key_type = double, value_type = short
2253 template<class key_type, class value_type>
2254 struct default_scan_by_key_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) <= 8)
2259  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
2260  && (sizeof(value_type) > 1))>>
2261  : scan_by_key_config_v2<64,
2262  14,
2263  ::rocprim::block_load_method::block_load_transpose,
2264  ::rocprim::block_store_method::block_store_transpose,
2265  block_scan_algorithm::using_warp_scan>
2266 {};
2267 
2268 // Based on key_type = double, value_type = int8_t
2269 template<class key_type, class value_type>
2270 struct default_scan_by_key_config<
2271  static_cast<unsigned int>(target_arch::gfx90a),
2272  key_type,
2273  value_type,
2274  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2275  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
2276  : scan_by_key_config_v2<64,
2277  14,
2278  ::rocprim::block_load_method::block_load_transpose,
2279  ::rocprim::block_store_method::block_store_transpose,
2280  block_scan_algorithm::reduce_then_scan>
2281 {};
2282 
2283 // Based on key_type = float, value_type = int64_t
2284 template<class key_type, class value_type>
2285 struct default_scan_by_key_config<
2286  static_cast<unsigned int>(target_arch::gfx90a),
2287  key_type,
2288  value_type,
2289  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2290  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
2291  && (sizeof(value_type) > 4))>>
2292  : scan_by_key_config_v2<128,
2293  10,
2294  ::rocprim::block_load_method::block_load_transpose,
2295  ::rocprim::block_store_method::block_store_transpose,
2296  block_scan_algorithm::using_warp_scan>
2297 {};
2298 
2299 // Based on key_type = float, value_type = int
2300 template<class key_type, class value_type>
2301 struct default_scan_by_key_config<
2302  static_cast<unsigned int>(target_arch::gfx90a),
2303  key_type,
2304  value_type,
2305  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2306  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
2307  && (sizeof(value_type) > 2))>>
2308  : scan_by_key_config_v2<256,
2309  12,
2310  ::rocprim::block_load_method::block_load_transpose,
2311  ::rocprim::block_store_method::block_store_transpose,
2312  block_scan_algorithm::reduce_then_scan>
2313 {};
2314 
2315 // Based on key_type = float, value_type = short
2316 template<class key_type, class value_type>
2317 struct default_scan_by_key_config<
2318  static_cast<unsigned int>(target_arch::gfx90a),
2319  key_type,
2320  value_type,
2321  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2322  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
2323  && (sizeof(value_type) > 1))>>
2324  : scan_by_key_config_v2<256,
2325  14,
2326  ::rocprim::block_load_method::block_load_transpose,
2327  ::rocprim::block_store_method::block_store_transpose,
2328  block_scan_algorithm::reduce_then_scan>
2329 {};
2330 
2331 // Based on key_type = float, value_type = int8_t
2332 template<class key_type, class value_type>
2333 struct default_scan_by_key_config<
2334  static_cast<unsigned int>(target_arch::gfx90a),
2335  key_type,
2336  value_type,
2337  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2338  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
2339  : scan_by_key_config_v2<128,
2340  14,
2341  ::rocprim::block_load_method::block_load_transpose,
2342  ::rocprim::block_store_method::block_store_transpose,
2343  block_scan_algorithm::using_warp_scan>
2344 {};
2345 
2346 // Based on key_type = rocprim::half, value_type = int64_t
2347 template<class key_type, class value_type>
2348 struct default_scan_by_key_config<
2349  static_cast<unsigned int>(target_arch::gfx90a),
2350  key_type,
2351  value_type,
2352  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2353  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
2354  : scan_by_key_config_v2<64,
2355  18,
2356  ::rocprim::block_load_method::block_load_transpose,
2357  ::rocprim::block_store_method::block_store_transpose,
2358  block_scan_algorithm::using_warp_scan>
2359 {};
2360 
2361 // Based on key_type = rocprim::half, value_type = int
2362 template<class key_type, class value_type>
2363 struct default_scan_by_key_config<
2364  static_cast<unsigned int>(target_arch::gfx90a),
2365  key_type,
2366  value_type,
2367  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2368  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
2369  : scan_by_key_config_v2<256,
2370  18,
2371  ::rocprim::block_load_method::block_load_transpose,
2372  ::rocprim::block_store_method::block_store_transpose,
2373  block_scan_algorithm::using_warp_scan>
2374 {};
2375 
2376 // Based on key_type = rocprim::half, value_type = short
2377 template<class key_type, class value_type>
2378 struct default_scan_by_key_config<
2379  static_cast<unsigned int>(target_arch::gfx90a),
2380  key_type,
2381  value_type,
2382  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2383  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
2384  : scan_by_key_config_v2<256,
2385  20,
2386  ::rocprim::block_load_method::block_load_transpose,
2387  ::rocprim::block_store_method::block_store_transpose,
2388  block_scan_algorithm::reduce_then_scan>
2389 {};
2390 
2391 // Based on key_type = rocprim::half, value_type = int8_t
2392 template<class key_type, class value_type>
2393 struct default_scan_by_key_config<
2394  static_cast<unsigned int>(target_arch::gfx90a),
2395  key_type,
2396  value_type,
2397  std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2398  && (sizeof(value_type) <= 1))>>
2399  : scan_by_key_config_v2<256,
2400  24,
2401  ::rocprim::block_load_method::block_load_transpose,
2402  ::rocprim::block_store_method::block_store_transpose,
2403  block_scan_algorithm::using_warp_scan>
2404 {};
2405 
2406 // Based on key_type = int64_t, value_type = int64_t
2407 template<class key_type, class value_type>
2408 struct default_scan_by_key_config<
2409  static_cast<unsigned int>(target_arch::gfx90a),
2410  key_type,
2411  value_type,
2412  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2413  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
2414  && (sizeof(value_type) > 4))>>
2415  : scan_by_key_config_v2<256,
2416  14,
2417  ::rocprim::block_load_method::block_load_transpose,
2418  ::rocprim::block_store_method::block_store_transpose,
2419  block_scan_algorithm::reduce_then_scan>
2420 {};
2421 
2422 // Based on key_type = int64_t, value_type = int
2423 template<class key_type, class value_type>
2424 struct default_scan_by_key_config<
2425  static_cast<unsigned int>(target_arch::gfx90a),
2426  key_type,
2427  value_type,
2428  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2429  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
2430  && (sizeof(value_type) > 2))>>
2431  : scan_by_key_config_v2<256,
2432  14,
2433  ::rocprim::block_load_method::block_load_transpose,
2434  ::rocprim::block_store_method::block_store_transpose,
2435  block_scan_algorithm::reduce_then_scan>
2436 {};
2437 
2438 // Based on key_type = int64_t, value_type = short
2439 template<class key_type, class value_type>
2440 struct default_scan_by_key_config<
2441  static_cast<unsigned int>(target_arch::gfx90a),
2442  key_type,
2443  value_type,
2444  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2445  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
2446  && (sizeof(value_type) > 1))>>
2447  : scan_by_key_config_v2<64,
2448  14,
2449  ::rocprim::block_load_method::block_load_transpose,
2450  ::rocprim::block_store_method::block_store_transpose,
2451  block_scan_algorithm::using_warp_scan>
2452 {};
2453 
2454 // Based on key_type = int64_t, value_type = int8_t
2455 template<class key_type, class value_type>
2456 struct default_scan_by_key_config<
2457  static_cast<unsigned int>(target_arch::gfx90a),
2458  key_type,
2459  value_type,
2460  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
2461  && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>>
2462  : scan_by_key_config_v2<256,
2463  14,
2464  ::rocprim::block_load_method::block_load_transpose,
2465  ::rocprim::block_store_method::block_store_transpose,
2466  block_scan_algorithm::using_warp_scan>
2467 {};
2468 
2469 // Based on key_type = int, value_type = int64_t
2470 template<class key_type, class value_type>
2471 struct default_scan_by_key_config<
2472  static_cast<unsigned int>(target_arch::gfx90a),
2473  key_type,
2474  value_type,
2475  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2476  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
2477  && (sizeof(value_type) > 4))>>
2478  : scan_by_key_config_v2<128,
2479  10,
2480  ::rocprim::block_load_method::block_load_transpose,
2481  ::rocprim::block_store_method::block_store_transpose,
2482  block_scan_algorithm::using_warp_scan>
2483 {};
2484 
2485 // Based on key_type = int, value_type = int
2486 template<class key_type, class value_type>
2487 struct default_scan_by_key_config<
2488  static_cast<unsigned int>(target_arch::gfx90a),
2489  key_type,
2490  value_type,
2491  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2492  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
2493  && (sizeof(value_type) > 2))>>
2494  : scan_by_key_config_v2<256,
2495  12,
2496  ::rocprim::block_load_method::block_load_transpose,
2497  ::rocprim::block_store_method::block_store_transpose,
2498  block_scan_algorithm::reduce_then_scan>
2499 {};
2500 
2501 // Based on key_type = int, value_type = short
2502 template<class key_type, class value_type>
2503 struct default_scan_by_key_config<
2504  static_cast<unsigned int>(target_arch::gfx90a),
2505  key_type,
2506  value_type,
2507  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2508  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
2509  && (sizeof(value_type) > 1))>>
2510  : scan_by_key_config_v2<256,
2511  14,
2512  ::rocprim::block_load_method::block_load_transpose,
2513  ::rocprim::block_store_method::block_store_transpose,
2514  block_scan_algorithm::reduce_then_scan>
2515 {};
2516 
2517 // Based on key_type = int, value_type = int8_t
2518 template<class key_type, class value_type>
2519 struct default_scan_by_key_config<
2520  static_cast<unsigned int>(target_arch::gfx90a),
2521  key_type,
2522  value_type,
2523  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
2524  && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>>
2525  : scan_by_key_config_v2<128,
2526  14,
2527  ::rocprim::block_load_method::block_load_transpose,
2528  ::rocprim::block_store_method::block_store_transpose,
2529  block_scan_algorithm::using_warp_scan>
2530 {};
2531 
2532 // Based on key_type = short, value_type = int64_t
2533 template<class key_type, class value_type>
2534 struct default_scan_by_key_config<
2535  static_cast<unsigned int>(target_arch::gfx90a),
2536  key_type,
2537  value_type,
2538  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2539  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
2540  && (sizeof(value_type) > 4))>>
2541  : scan_by_key_config_v2<64,
2542  18,
2543  ::rocprim::block_load_method::block_load_transpose,
2544  ::rocprim::block_store_method::block_store_transpose,
2545  block_scan_algorithm::reduce_then_scan>
2546 {};
2547 
2548 // Based on key_type = short, value_type = int
2549 template<class key_type, class value_type>
2550 struct default_scan_by_key_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) && (sizeof(value_type) <= 4)
2556  && (sizeof(value_type) > 2))>>
2557  : scan_by_key_config_v2<256,
2558  18,
2559  ::rocprim::block_load_method::block_load_transpose,
2560  ::rocprim::block_store_method::block_store_transpose,
2561  block_scan_algorithm::using_warp_scan>
2562 {};
2563 
2564 // Based on key_type = short, value_type = short
2565 template<class key_type, class value_type>
2566 struct default_scan_by_key_config<
2567  static_cast<unsigned int>(target_arch::gfx90a),
2568  key_type,
2569  value_type,
2570  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2571  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
2572  && (sizeof(value_type) > 1))>>
2573  : scan_by_key_config_v2<256,
2574  20,
2575  ::rocprim::block_load_method::block_load_transpose,
2576  ::rocprim::block_store_method::block_store_transpose,
2577  block_scan_algorithm::reduce_then_scan>
2578 {};
2579 
2580 // Based on key_type = short, value_type = int8_t
2581 template<class key_type, class value_type>
2582 struct default_scan_by_key_config<
2583  static_cast<unsigned int>(target_arch::gfx90a),
2584  key_type,
2585  value_type,
2586  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
2587  && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1))>>
2588  : scan_by_key_config_v2<256,
2589  24,
2590  ::rocprim::block_load_method::block_load_transpose,
2591  ::rocprim::block_store_method::block_store_transpose,
2592  block_scan_algorithm::using_warp_scan>
2593 {};
2594 
2595 // Based on key_type = int8_t, value_type = int64_t
2596 template<class key_type, class value_type>
2597 struct default_scan_by_key_config<
2598  static_cast<unsigned int>(target_arch::gfx90a),
2599  key_type,
2600  value_type,
2601  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2602  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
2603  : scan_by_key_config_v2<64,
2604  18,
2605  ::rocprim::block_load_method::block_load_transpose,
2606  ::rocprim::block_store_method::block_store_transpose,
2607  block_scan_algorithm::using_warp_scan>
2608 {};
2609 
2610 // Based on key_type = int8_t, value_type = int
2611 template<class key_type, class value_type>
2612 struct default_scan_by_key_config<
2613  static_cast<unsigned int>(target_arch::gfx90a),
2614  key_type,
2615  value_type,
2616  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2617  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
2618  : scan_by_key_config_v2<128,
2619  18,
2620  ::rocprim::block_load_method::block_load_transpose,
2621  ::rocprim::block_store_method::block_store_transpose,
2622  block_scan_algorithm::using_warp_scan>
2623 {};
2624 
2625 // Based on key_type = int8_t, value_type = short
2626 template<class key_type, class value_type>
2627 struct default_scan_by_key_config<
2628  static_cast<unsigned int>(target_arch::gfx90a),
2629  key_type,
2630  value_type,
2631  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2632  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
2633  : scan_by_key_config_v2<256,
2634  20,
2635  ::rocprim::block_load_method::block_load_transpose,
2636  ::rocprim::block_store_method::block_store_transpose,
2637  block_scan_algorithm::reduce_then_scan>
2638 {};
2639 
2640 // Based on key_type = int8_t, value_type = int8_t
2641 template<class key_type, class value_type>
2642 struct default_scan_by_key_config<
2643  static_cast<unsigned int>(target_arch::gfx90a),
2644  key_type,
2645  value_type,
2646  std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
2647  && (sizeof(value_type) <= 1))>>
2648  : scan_by_key_config_v2<256,
2649  24,
2650  ::rocprim::block_load_method::block_load_transpose,
2651  ::rocprim::block_store_method::block_store_transpose,
2652  block_scan_algorithm::reduce_then_scan>
2653 {};
2654 
2655 } // end namespace detail
2656 
2657 END_ROCPRIM_NAMESPACE
2658 
2660 // end of group primitivesmodule_deviceconfigs
2661 
2662 #endif // ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_SCAN_BY_KEY_HPP_
scan_by_key_config_v2
Configuration of device-level scan-by-key operation.
Definition: device_config_helper.hpp:417
detail::default_scan_by_key_config_base
Definition: device_config_helper.hpp:513
detail
Deprecated: Configuration of device-level scan primitives.
Definition: block_histogram.hpp:62
detail::default_scan_by_key_config
Definition: device_scan_by_key.hpp:42
Generated by   doxygen 1.8.13