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