rocPRIM
device_merge_config.hpp
1 // Copyright (c) 2018-2022 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_DEVICE_MERGE_CONFIG_HPP_
22 #define ROCPRIM_DEVICE_DEVICE_MERGE_CONFIG_HPP_
23 
24 #include <type_traits>
25 
26 #include "../config.hpp"
27 #include "../detail/various.hpp"
28 #include "../functional.hpp"
29 
30 #include "config_types.hpp"
31 
34 
35 BEGIN_ROCPRIM_NAMESPACE
36 
38 template<unsigned int BlockSize, unsigned int ItemsPerThread>
40 
41 namespace detail
42 {
43 
44 template<class Key, class Value>
46 {
47  static constexpr unsigned int item_scale =
48  ::rocprim::detail::ceiling_div<unsigned int>(::rocprim::max(sizeof(Key), sizeof(Value)), sizeof(int));
49 
50  // TODO Tune when merge-by-key is ready
51  using type = merge_config<256, ::rocprim::max(1u, 10u / item_scale)>;
52 };
53 
54 template<class Key>
56 {
57  static constexpr unsigned int item_scale =
58  ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Key), sizeof(int));
59 
60  using type = select_type<
64  merge_config<256, ::rocprim::max(1u, 10u / item_scale)>
65  >;
66 };
67 
68 template<class Key, class Value>
70 {
71  static constexpr unsigned int item_scale =
72  ::rocprim::detail::ceiling_div<unsigned int>(::rocprim::max(sizeof(Key), sizeof(Value)), sizeof(int));
73 
74  // TODO Tune when merge-by-key is ready
75  using type = merge_config<256, ::rocprim::max(1u, 10u / item_scale)>;
76 };
77 
78 template<class Key>
80 {
81  static constexpr unsigned int item_scale =
82  ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Key), sizeof(int));
83 
84  using type = select_type<
88  merge_config<256, ::rocprim::max(1u, 10u / item_scale)>
89  >;
90 };
91 
92 // TODO: We need to update these parameters
93 template<class Key, class Value>
95 {
96  static constexpr unsigned int item_scale =
97  ::rocprim::detail::ceiling_div<unsigned int>(::rocprim::max(sizeof(Key), sizeof(Value)), sizeof(int));
98 
99  // TODO Tune when merge-by-key is ready
100  using type = merge_config<256, ::rocprim::max(1u, 10u / item_scale)>;
101 };
102 
103 template<class Key>
105 {
106  static constexpr unsigned int item_scale =
107  ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Key), sizeof(int));
108 
109  using type = select_type<
113  merge_config<256, ::rocprim::max(1u, 10u / item_scale)>
114  >;
115 };
116 
117 // TODO: We need to update these parameters
118 template<class Key, class Value>
120 {
121  static constexpr unsigned int item_scale =
122  ::rocprim::detail::ceiling_div<unsigned int>(::rocprim::max(sizeof(Key), sizeof(Value)), sizeof(int));
123 
124  // TODO Tune when merge-by-key is ready
125  using type = merge_config<256, ::rocprim::max(1u, 10u / item_scale)>;
126 };
127 
128 template<class Key>
130 {
131  static constexpr unsigned int item_scale =
132  ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Key), sizeof(int));
133 
134  using type = select_type<
138  merge_config<256, ::rocprim::max(1u, 10u / item_scale)>
139  >;
140 };
141 
142 template<unsigned int TargetArch, class Key, class Value>
144  : select_arch<
145  TargetArch,
146  select_arch_case<803, merge_config_803<Key, Value>>,
147  select_arch_case<900, merge_config_900<Key, Value>>,
148  select_arch_case<ROCPRIM_ARCH_90a, merge_config_90a<Key, Value>>,
149  select_arch_case<1030, merge_config_1030<Key, Value>>,
150  merge_config_900<Key, Value>
151  > { };
152 
153 } // end namespace detail
154 
155 END_ROCPRIM_NAMESPACE
156 
158 // end of group primitivesmodule_deviceconfigs
159 
160 #endif // ROCPRIM_DEVICE_DEVICE_MERGE_CONFIG_HPP_
Empty type used as a placeholder, usually used to flag that given template parameter should not be us...
Definition: types.hpp:135
Definition: device_merge_config.hpp:119
ROCPRIM_HOST_DEVICE constexpr T max(const T &a, const T &b)
Returns the maximum of its arguments.
Definition: functional.hpp:55
Definition: device_merge_config.hpp:45
Definition: various.hpp:236
Definition: device_merge_config.hpp:143
Definition: device_merge_config.hpp:69
Deprecated: Configuration of device-level scan primitives.
Definition: block_histogram.hpp:62
Definition: config_types.hpp:140
Definition: device_merge_config.hpp:94
Configuration of particular kernels launched by device-level operation.
Definition: config_types.hpp:84