rocPRIM
device_scan.hpp
1 // Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved.
2 //
3 // Permission is hereby granted, free of charge, to any person obtaining a copy
4 // of this software and associated documentation files (the "Software"), to deal
5 // in the Software without restriction, including without limitation the rights
6 // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7 // copies of the Software, and to permit persons to whom the Software is
8 // furnished to do so, subject to the following conditions:
9 //
10 // The above copyright notice and this permission notice shall be included in
11 // all copies or substantial portions of the Software.
12 //
13 // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
14 // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
15 // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
16 // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
17 // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
18 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
19 // THE SOFTWARE.
20 
21 #ifndef ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_SCAN_HPP_
22 #define ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_SCAN_HPP_
23 
24 #include "../../../type_traits.hpp"
25 #include "../device_config_helper.hpp"
26 #include <type_traits>
27 
28 /* DO NOT EDIT THIS FILE
29  * This file is automatically generated by `/scripts/autotune/create_optimization.py`.
30  * so most likely you want to edit rocprim/device/device_(algo)_config.hpp
31  */
32 
35 
36 BEGIN_ROCPRIM_NAMESPACE
37 
38 namespace detail
39 {
40 
41 template<unsigned int arch, class value_type, class enable = void>
42 struct default_scan_config : default_scan_config_base<value_type>
43 {};
44 
45 // Based on value_type = double
46 template<class value_type>
47 struct default_scan_config<
48  static_cast<unsigned int>(target_arch::gfx908),
49  value_type,
50  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
51  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
52  : scan_config_v2<256,
53  6,
54  ::rocprim::block_load_method::block_load_transpose,
55  ::rocprim::block_store_method::block_store_transpose,
56  block_scan_algorithm::reduce_then_scan>
57 {};
58 
59 // Based on value_type = float
60 template<class value_type>
61 struct default_scan_config<
62  static_cast<unsigned int>(target_arch::gfx908),
63  value_type,
64  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
65  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
66  : scan_config_v2<256,
67  14,
68  ::rocprim::block_load_method::block_load_transpose,
69  ::rocprim::block_store_method::block_store_transpose,
70  block_scan_algorithm::reduce_then_scan>
71 {};
72 
73 // Based on value_type = rocprim::half
74 template<class value_type>
75 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx908),
76  value_type,
77  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
78  && (sizeof(value_type) <= 2))>>
79  : scan_config_v2<256,
80  18,
81  ::rocprim::block_load_method::block_load_transpose,
82  ::rocprim::block_store_method::block_store_transpose,
83  block_scan_algorithm::reduce_then_scan>
84 {};
85 
86 // Based on value_type = int64_t
87 template<class value_type>
88 struct default_scan_config<
89  static_cast<unsigned int>(target_arch::gfx908),
90  value_type,
91  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
92  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
93  : scan_config_v2<256,
94  6,
95  ::rocprim::block_load_method::block_load_transpose,
96  ::rocprim::block_store_method::block_store_transpose,
97  block_scan_algorithm::using_warp_scan>
98 {};
99 
100 // Based on value_type = int
101 template<class value_type>
102 struct default_scan_config<
103  static_cast<unsigned int>(target_arch::gfx908),
104  value_type,
105  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
106  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
107  : scan_config_v2<256,
108  14,
109  ::rocprim::block_load_method::block_load_transpose,
110  ::rocprim::block_store_method::block_store_transpose,
111  block_scan_algorithm::using_warp_scan>
112 {};
113 
114 // Based on value_type = short
115 template<class value_type>
116 struct default_scan_config<
117  static_cast<unsigned int>(target_arch::gfx908),
118  value_type,
119  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
120  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
121  : scan_config_v2<128,
122  14,
123  ::rocprim::block_load_method::block_load_transpose,
124  ::rocprim::block_store_method::block_store_transpose,
125  block_scan_algorithm::using_warp_scan>
126 {};
127 
128 // Based on value_type = int8_t
129 template<class value_type>
130 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx908),
131  value_type,
132  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
133  && (sizeof(value_type) <= 1))>>
134  : scan_config_v2<256,
135  24,
136  ::rocprim::block_load_method::block_load_transpose,
137  ::rocprim::block_store_method::block_store_transpose,
138  block_scan_algorithm::using_warp_scan>
139 {};
140 
141 // Based on value_type = double
142 template<class value_type>
143 struct default_scan_config<
144  static_cast<unsigned int>(target_arch::gfx900),
145  value_type,
146  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
147  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
148  : scan_config_v2<256,
149  15,
150  ::rocprim::block_load_method::block_load_transpose,
151  ::rocprim::block_store_method::block_store_transpose,
152  block_scan_algorithm::reduce_then_scan>
153 {};
154 
155 // Based on value_type = float
156 template<class value_type>
157 struct default_scan_config<
158  static_cast<unsigned int>(target_arch::gfx900),
159  value_type,
160  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
161  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
162  : scan_config_v2<256,
163  10,
164  ::rocprim::block_load_method::block_load_transpose,
165  ::rocprim::block_store_method::block_store_transpose,
166  block_scan_algorithm::using_warp_scan>
167 {};
168 
169 // Based on value_type = rocprim::half
170 template<class value_type>
171 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx900),
172  value_type,
173  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
174  && (sizeof(value_type) <= 2))>>
175  : scan_config_v2<256,
176  24,
177  ::rocprim::block_load_method::block_load_transpose,
178  ::rocprim::block_store_method::block_store_transpose,
179  block_scan_algorithm::using_warp_scan>
180 {};
181 
182 // Based on value_type = int64_t
183 template<class value_type>
184 struct default_scan_config<
185  static_cast<unsigned int>(target_arch::gfx900),
186  value_type,
187  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
188  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
189  : scan_config_v2<256,
190  10,
191  ::rocprim::block_load_method::block_load_transpose,
192  ::rocprim::block_store_method::block_store_transpose,
193  block_scan_algorithm::using_warp_scan>
194 {};
195 
196 // Based on value_type = int
197 template<class value_type>
198 struct default_scan_config<
199  static_cast<unsigned int>(target_arch::gfx900),
200  value_type,
201  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
202  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
203  : scan_config_v2<256,
204  10,
205  ::rocprim::block_load_method::block_load_transpose,
206  ::rocprim::block_store_method::block_store_transpose,
207  block_scan_algorithm::using_warp_scan>
208 {};
209 
210 // Based on value_type = short
211 template<class value_type>
212 struct default_scan_config<
213  static_cast<unsigned int>(target_arch::gfx900),
214  value_type,
215  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
216  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
217  : scan_config_v2<256,
218  24,
219  ::rocprim::block_load_method::block_load_transpose,
220  ::rocprim::block_store_method::block_store_transpose,
221  block_scan_algorithm::using_warp_scan>
222 {};
223 
224 // Based on value_type = int8_t
225 template<class value_type>
226 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx900),
227  value_type,
228  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
229  && (sizeof(value_type) <= 1))>>
230  : scan_config_v2<256,
231  24,
232  ::rocprim::block_load_method::block_load_transpose,
233  ::rocprim::block_store_method::block_store_transpose,
234  block_scan_algorithm::using_warp_scan>
235 {};
236 
237 // Based on value_type = double
238 template<class value_type>
239 struct default_scan_config<
240  static_cast<unsigned int>(target_arch::gfx906),
241  value_type,
242  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
243  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
244  : scan_config_v2<128,
245  10,
246  ::rocprim::block_load_method::block_load_transpose,
247  ::rocprim::block_store_method::block_store_transpose,
248  block_scan_algorithm::reduce_then_scan>
249 {};
250 
251 // Based on value_type = float
252 template<class value_type>
253 struct default_scan_config<
254  static_cast<unsigned int>(target_arch::gfx906),
255  value_type,
256  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
257  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
258  : scan_config_v2<256,
259  15,
260  ::rocprim::block_load_method::block_load_transpose,
261  ::rocprim::block_store_method::block_store_transpose,
262  block_scan_algorithm::using_warp_scan>
263 {};
264 
265 // Based on value_type = rocprim::half
266 template<class value_type>
267 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx906),
268  value_type,
269  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
270  && (sizeof(value_type) <= 2))>>
271  : scan_config_v2<256,
272  24,
273  ::rocprim::block_load_method::block_load_transpose,
274  ::rocprim::block_store_method::block_store_transpose,
275  block_scan_algorithm::reduce_then_scan>
276 {};
277 
278 // Based on value_type = int64_t
279 template<class value_type>
280 struct default_scan_config<
281  static_cast<unsigned int>(target_arch::gfx906),
282  value_type,
283  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
284  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
285  : scan_config_v2<64,
286  14,
287  ::rocprim::block_load_method::block_load_transpose,
288  ::rocprim::block_store_method::block_store_transpose,
289  block_scan_algorithm::using_warp_scan>
290 {};
291 
292 // Based on value_type = int
293 template<class value_type>
294 struct default_scan_config<
295  static_cast<unsigned int>(target_arch::gfx906),
296  value_type,
297  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
298  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
299  : scan_config_v2<256,
300  15,
301  ::rocprim::block_load_method::block_load_transpose,
302  ::rocprim::block_store_method::block_store_transpose,
303  block_scan_algorithm::reduce_then_scan>
304 {};
305 
306 // Based on value_type = short
307 template<class value_type>
308 struct default_scan_config<
309  static_cast<unsigned int>(target_arch::gfx906),
310  value_type,
311  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
312  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
313  : scan_config_v2<256,
314  24,
315  ::rocprim::block_load_method::block_load_transpose,
316  ::rocprim::block_store_method::block_store_transpose,
317  block_scan_algorithm::reduce_then_scan>
318 {};
319 
320 // Based on value_type = int8_t
321 template<class value_type>
322 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx906),
323  value_type,
324  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
325  && (sizeof(value_type) <= 1))>>
326  : scan_config_v2<256,
327  24,
328  ::rocprim::block_load_method::block_load_transpose,
329  ::rocprim::block_store_method::block_store_transpose,
330  block_scan_algorithm::using_warp_scan>
331 {};
332 
333 // Based on value_type = double
334 template<class value_type>
335 struct default_scan_config<
336  static_cast<unsigned int>(target_arch::gfx1030),
337  value_type,
338  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
339  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
340  : scan_config_v2<128,
341  14,
342  ::rocprim::block_load_method::block_load_transpose,
343  ::rocprim::block_store_method::block_store_transpose,
344  block_scan_algorithm::using_warp_scan>
345 {};
346 
347 // Based on value_type = float
348 template<class value_type>
349 struct default_scan_config<
350  static_cast<unsigned int>(target_arch::gfx1030),
351  value_type,
352  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
353  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
354  : scan_config_v2<64,
355  18,
356  ::rocprim::block_load_method::block_load_transpose,
357  ::rocprim::block_store_method::block_store_transpose,
358  block_scan_algorithm::using_warp_scan>
359 {};
360 
361 // Based on value_type = rocprim::half
362 template<class value_type>
363 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx1030),
364  value_type,
365  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
366  && (sizeof(value_type) <= 2))>>
367  : scan_config_v2<256,
368  22,
369  ::rocprim::block_load_method::block_load_transpose,
370  ::rocprim::block_store_method::block_store_transpose,
371  block_scan_algorithm::using_warp_scan>
372 {};
373 
374 // Based on value_type = int64_t
375 template<class value_type>
376 struct default_scan_config<
377  static_cast<unsigned int>(target_arch::gfx1030),
378  value_type,
379  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
380  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
381  : scan_config_v2<256,
382  9,
383  ::rocprim::block_load_method::block_load_transpose,
384  ::rocprim::block_store_method::block_store_transpose,
385  block_scan_algorithm::using_warp_scan>
386 {};
387 
388 // Based on value_type = int
389 template<class value_type>
390 struct default_scan_config<
391  static_cast<unsigned int>(target_arch::gfx1030),
392  value_type,
393  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
394  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
395  : scan_config_v2<64,
396  18,
397  ::rocprim::block_load_method::block_load_transpose,
398  ::rocprim::block_store_method::block_store_transpose,
399  block_scan_algorithm::using_warp_scan>
400 {};
401 
402 // Based on value_type = short
403 template<class value_type>
404 struct default_scan_config<
405  static_cast<unsigned int>(target_arch::gfx1030),
406  value_type,
407  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
408  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
409  : scan_config_v2<256,
410  22,
411  ::rocprim::block_load_method::block_load_transpose,
412  ::rocprim::block_store_method::block_store_transpose,
413  block_scan_algorithm::using_warp_scan>
414 {};
415 
416 // Based on value_type = int8_t
417 template<class value_type>
418 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx1030),
419  value_type,
420  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
421  && (sizeof(value_type) <= 1))>>
422  : scan_config_v2<256,
423  24,
424  ::rocprim::block_load_method::block_load_transpose,
425  ::rocprim::block_store_method::block_store_transpose,
426  block_scan_algorithm::using_warp_scan>
427 {};
428 
429 // Based on value_type = double
430 template<class value_type>
431 struct default_scan_config<
432  static_cast<unsigned int>(target_arch::unknown),
433  value_type,
434  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
435  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
436  : scan_config_v2<256,
437  6,
438  ::rocprim::block_load_method::block_load_transpose,
439  ::rocprim::block_store_method::block_store_transpose,
440  block_scan_algorithm::reduce_then_scan>
441 {};
442 
443 // Based on value_type = float
444 template<class value_type>
445 struct default_scan_config<
446  static_cast<unsigned int>(target_arch::unknown),
447  value_type,
448  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
449  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
450  : scan_config_v2<256,
451  14,
452  ::rocprim::block_load_method::block_load_transpose,
453  ::rocprim::block_store_method::block_store_transpose,
454  block_scan_algorithm::reduce_then_scan>
455 {};
456 
457 // Based on value_type = rocprim::half
458 template<class value_type>
459 struct default_scan_config<static_cast<unsigned int>(target_arch::unknown),
460  value_type,
461  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
462  && (sizeof(value_type) <= 2))>>
463  : scan_config_v2<256,
464  18,
465  ::rocprim::block_load_method::block_load_transpose,
466  ::rocprim::block_store_method::block_store_transpose,
467  block_scan_algorithm::reduce_then_scan>
468 {};
469 
470 // Based on value_type = int64_t
471 template<class value_type>
472 struct default_scan_config<
473  static_cast<unsigned int>(target_arch::unknown),
474  value_type,
475  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
476  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
477  : scan_config_v2<256,
478  6,
479  ::rocprim::block_load_method::block_load_transpose,
480  ::rocprim::block_store_method::block_store_transpose,
481  block_scan_algorithm::using_warp_scan>
482 {};
483 
484 // Based on value_type = int
485 template<class value_type>
486 struct default_scan_config<
487  static_cast<unsigned int>(target_arch::unknown),
488  value_type,
489  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
490  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
491  : scan_config_v2<256,
492  14,
493  ::rocprim::block_load_method::block_load_transpose,
494  ::rocprim::block_store_method::block_store_transpose,
495  block_scan_algorithm::using_warp_scan>
496 {};
497 
498 // Based on value_type = short
499 template<class value_type>
500 struct default_scan_config<
501  static_cast<unsigned int>(target_arch::unknown),
502  value_type,
503  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
504  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
505  : scan_config_v2<128,
506  14,
507  ::rocprim::block_load_method::block_load_transpose,
508  ::rocprim::block_store_method::block_store_transpose,
509  block_scan_algorithm::using_warp_scan>
510 {};
511 
512 // Based on value_type = int8_t
513 template<class value_type>
514 struct default_scan_config<static_cast<unsigned int>(target_arch::unknown),
515  value_type,
516  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
517  && (sizeof(value_type) <= 1))>>
518  : scan_config_v2<256,
519  24,
520  ::rocprim::block_load_method::block_load_transpose,
521  ::rocprim::block_store_method::block_store_transpose,
522  block_scan_algorithm::using_warp_scan>
523 {};
524 
525 // Based on value_type = double
526 template<class value_type>
527 struct default_scan_config<
528  static_cast<unsigned int>(target_arch::gfx90a),
529  value_type,
530  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
531  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
532  : scan_config_v2<256,
533  6,
534  ::rocprim::block_load_method::block_load_transpose,
535  ::rocprim::block_store_method::block_store_transpose,
536  block_scan_algorithm::reduce_then_scan>
537 {};
538 
539 // Based on value_type = float
540 template<class value_type>
541 struct default_scan_config<
542  static_cast<unsigned int>(target_arch::gfx90a),
543  value_type,
544  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
545  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
546  : scan_config_v2<256,
547  14,
548  ::rocprim::block_load_method::block_load_transpose,
549  ::rocprim::block_store_method::block_store_transpose,
550  block_scan_algorithm::reduce_then_scan>
551 {};
552 
553 // Based on value_type = rocprim::half
554 template<class value_type>
555 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx90a),
556  value_type,
557  std::enable_if_t<(bool(rocprim::is_floating_point<value_type>::value)
558  && (sizeof(value_type) <= 2))>>
559  : scan_config_v2<256,
560  18,
561  ::rocprim::block_load_method::block_load_transpose,
562  ::rocprim::block_store_method::block_store_transpose,
563  block_scan_algorithm::reduce_then_scan>
564 {};
565 
566 // Based on value_type = int64_t
567 template<class value_type>
568 struct default_scan_config<
569  static_cast<unsigned int>(target_arch::gfx90a),
570  value_type,
571  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
572  && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
573  : scan_config_v2<256,
574  6,
575  ::rocprim::block_load_method::block_load_transpose,
576  ::rocprim::block_store_method::block_store_transpose,
577  block_scan_algorithm::using_warp_scan>
578 {};
579 
580 // Based on value_type = int
581 template<class value_type>
582 struct default_scan_config<
583  static_cast<unsigned int>(target_arch::gfx90a),
584  value_type,
585  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
586  && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
587  : scan_config_v2<256,
588  14,
589  ::rocprim::block_load_method::block_load_transpose,
590  ::rocprim::block_store_method::block_store_transpose,
591  block_scan_algorithm::using_warp_scan>
592 {};
593 
594 // Based on value_type = short
595 template<class value_type>
596 struct default_scan_config<
597  static_cast<unsigned int>(target_arch::gfx90a),
598  value_type,
599  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
600  && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
601  : scan_config_v2<128,
602  14,
603  ::rocprim::block_load_method::block_load_transpose,
604  ::rocprim::block_store_method::block_store_transpose,
605  block_scan_algorithm::using_warp_scan>
606 {};
607 
608 // Based on value_type = int8_t
609 template<class value_type>
610 struct default_scan_config<static_cast<unsigned int>(target_arch::gfx90a),
611  value_type,
612  std::enable_if_t<(!bool(rocprim::is_floating_point<value_type>::value)
613  && (sizeof(value_type) <= 1))>>
614  : scan_config_v2<256,
615  24,
616  ::rocprim::block_load_method::block_load_transpose,
617  ::rocprim::block_store_method::block_store_transpose,
618  block_scan_algorithm::using_warp_scan>
619 {};
620 
621 } // end namespace detail
622 
623 END_ROCPRIM_NAMESPACE
624 
626 // end of group primitivesmodule_deviceconfigs
627 
628 #endif // ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_SCAN_HPP_
detail
Deprecated: Configuration of device-level scan primitives.
Definition: block_histogram.hpp:62
detail::default_scan_config_base
Definition: device_config_helper.hpp:388
detail::default_scan_config
Definition: device_scan.hpp:42
scan_config_v2
Configuration of device-level scan primitives.
Definition: device_config_helper.hpp:294
Generated by   doxygen 1.8.13