21 #ifndef ROCPRIM_DEVICE_DEVICE_SELECT_CONFIG_HPP_ 22 #define ROCPRIM_DEVICE_DEVICE_SELECT_CONFIG_HPP_ 24 #include <type_traits> 26 #include "../config.hpp" 27 #include "../detail/various.hpp" 28 #include "../functional.hpp" 30 #include "../block/block_load.hpp" 31 #include "../block/block_scan.hpp" 33 #include "config_types.hpp" 38 BEGIN_ROCPRIM_NAMESPACE
50 unsigned int BlockSize,
51 unsigned int ItemsPerThread,
56 unsigned int SizeLimit = ROCPRIM_GRID_SIZE_LIMIT
82 static constexpr
unsigned int item_scale =
83 ::rocprim::detail::ceiling_div<unsigned int>(
sizeof(Key),
sizeof(
int));
88 ::rocprim::block_load_method::block_load_transpose,
89 ::rocprim::block_load_method::block_load_transpose,
90 ::rocprim::block_load_method::block_load_transpose,
91 ::rocprim::block_scan_algorithm::using_warp_scan
98 static constexpr
unsigned int item_scale =
99 ::rocprim::detail::ceiling_div<unsigned int>(
sizeof(Key),
sizeof(
int));
104 ::rocprim::block_load_method::block_load_transpose,
105 ::rocprim::block_load_method::block_load_transpose,
106 ::rocprim::block_load_method::block_load_transpose,
107 ::rocprim::block_scan_algorithm::using_warp_scan
111 template<
class Value>
114 static constexpr
unsigned int item_scale =
115 ::rocprim::detail::ceiling_div<unsigned int>(
sizeof(Value),
sizeof(
int));
120 ::rocprim::block_load_method::block_load_transpose,
121 ::rocprim::block_load_method::block_load_transpose,
122 ::rocprim::block_load_method::block_load_transpose,
123 ::rocprim::block_scan_algorithm::using_warp_scan
127 template<
class Value>
130 static constexpr
unsigned int item_scale =
131 ::rocprim::detail::ceiling_div<unsigned int>(
sizeof(Value),
sizeof(
int));
136 ::rocprim::block_load_method::block_load_transpose,
137 ::rocprim::block_load_method::block_load_transpose,
138 ::rocprim::block_load_method::block_load_transpose,
139 ::rocprim::block_scan_algorithm::using_warp_scan
144 template<
unsigned int TargetArch,
class Key,
class >
148 select_arch_case<803, select_config_803<Key>>,
149 select_arch_case<900, select_config_900<Key>>,
150 select_arch_case<ROCPRIM_ARCH_90a, select_config_90a<Key>>,
151 select_arch_case<1030, select_config_1030<Key>>,
152 select_config_803<Key>
157 END_ROCPRIM_NAMESPACE
162 #endif // ROCPRIM_DEVICE_DEVICE_SELECT_CONFIG_HPP_ ROCPRIM_HOST_DEVICE constexpr T max(const T &a, const T &b)
Returns the maximum of its arguments.
Definition: functional.hpp:55
Definition: device_select_config.hpp:145
Definition: config_types.hpp:109
static constexpr block_load_method value_block_load_method
Method for loading input values.
Definition: device_select_config.hpp:67
Configuration of device-level select operation.
Definition: device_select_config.hpp:58
static constexpr unsigned int size_limit
Limit on the number of items for a single select kernel launch.
Definition: device_select_config.hpp:73
static constexpr unsigned int block_size
Number of threads in a block.
Definition: device_select_config.hpp:61
Deprecated: Configuration of device-level scan primitives.
Definition: block_histogram.hpp:62
static constexpr block_load_method key_block_load_method
Method for loading input keys.
Definition: device_select_config.hpp:65
block_load_method
block_load_method enumerates the methods available to load data from continuous memory into a blocked...
Definition: block_load.hpp:41
Definition: device_select_config.hpp:112
Definition: config_types.hpp:140
static constexpr unsigned int items_per_thread
Number of items processed by each thread.
Definition: device_select_config.hpp:63
Definition: device_select_config.hpp:80
static constexpr block_scan_algorithm block_scan_method
Algorithm for block scan.
Definition: device_select_config.hpp:71
Definition: device_select_config.hpp:96
static constexpr block_load_method flag_block_load_method
Method for loading flag values.
Definition: device_select_config.hpp:69
Definition: device_select_config.hpp:128
block_scan_algorithm
Available algorithms for block_scan primitive.
Definition: block_scan.hpp:41