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