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