diff --git a/rocprim/include/rocprim/device/config_types.hpp b/rocprim/include/rocprim/device/config_types.hpp index a27993345..2a7c7115b 100644 --- a/rocprim/include/rocprim/device/config_types.hpp +++ b/rocprim/include/rocprim/device/config_types.hpp @@ -173,6 +173,8 @@ enum class target_arch : unsigned int gfx1030 = 1030, gfx1100 = 1100, gfx1102 = 1102, + gfx1200 = 1200, + gfx1201 = 1201, unknown = std::numeric_limits::max(), }; #endif // DOXYGEN_SHOULD_SKIP_THIS @@ -212,7 +214,9 @@ constexpr target_arch get_target_arch_from_name(const char* const arch_name, con "gfx942", "gfx1030", "gfx1100", - "gfx1102"}; + "gfx1102", + "gfx1200", + "gfx1201"}; constexpr target_arch target_architectures[] = { target_arch::gfx803, target_arch::gfx900, @@ -223,6 +227,8 @@ constexpr target_arch get_target_arch_from_name(const char* const arch_name, con target_arch::gfx1030, target_arch::gfx1100, target_arch::gfx1102, + target_arch::gfx1200, + target_arch::gfx1201, }; static_assert(sizeof(target_names) / sizeof(target_names[0]) == sizeof(target_architectures) / sizeof(target_architectures[0]), @@ -283,6 +289,10 @@ auto dispatch_target_arch(const target_arch target_arch) return Config::template architecture_config::params; case target_arch::gfx1102: return Config::template architecture_config::params; + case target_arch::gfx1200: + return Config::template architecture_config::params; + case target_arch::gfx1201: + return Config::template architecture_config::params; case target_arch::invalid: assert(false && "Invalid target architecture selected at runtime."); } diff --git a/rocprim/include/rocprim/device/detail/config/device_adjacent_find.hpp b/rocprim/include/rocprim/device/detail/config/device_adjacent_find.hpp index 2b3609c60..16951744f 100644 --- a/rocprim/include/rocprim/device/detail/config/device_adjacent_find.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_adjacent_find.hpp @@ -201,6 +201,142 @@ struct default_adjacent_find_config< : adjacent_find_config<1024, 4> {}; +// Based on input_type = double +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1200), + input_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 8) && (sizeof(input_type) > 4))>> + : adjacent_find_config<64, 2> +{}; + +// Based on input_type = float +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1200), + input_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 4) && (sizeof(input_type) > 2))>> + : adjacent_find_config<128, 4> +{}; + +// Based on input_type = rocprim::half +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1200), + input_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 2))>> : adjacent_find_config<64, 64> +{}; + +// Based on input_type = int64_t +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1200), + input_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 8) && (sizeof(input_type) > 4))>> + : adjacent_find_config<64, 4> +{}; + +// Based on input_type = int +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1200), + input_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 4) && (sizeof(input_type) > 2))>> + : adjacent_find_config<128, 32> +{}; + +// Based on input_type = short +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1200), + input_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 2) && (sizeof(input_type) > 1))>> + : adjacent_find_config<256, 32> +{}; + +// Based on input_type = int8_t +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1200), + input_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 1))>> : adjacent_find_config<256, 64> +{}; + +// Based on input_type = double +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1201), + input_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 8) && (sizeof(input_type) > 4))>> + : adjacent_find_config<64, 2> +{}; + +// Based on input_type = float +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1201), + input_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 4) && (sizeof(input_type) > 2))>> + : adjacent_find_config<64, 32> +{}; + +// Based on input_type = rocprim::half +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1201), + input_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 2))>> : adjacent_find_config<256, 16> +{}; + +// Based on input_type = int64_t +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1201), + input_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 8) && (sizeof(input_type) > 4))>> + : adjacent_find_config<512, 2> +{}; + +// Based on input_type = int +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1201), + input_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 4) && (sizeof(input_type) > 2))>> + : adjacent_find_config<64, 16> +{}; + +// Based on input_type = short +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1201), + input_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 2) && (sizeof(input_type) > 1))>> + : adjacent_find_config<256, 2> +{}; + +// Based on input_type = int8_t +template +struct default_adjacent_find_config< + static_cast(target_arch::gfx1201), + input_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(input_type) <= 1))>> : adjacent_find_config<128, 64> +{}; + // Based on input_type = double template struct default_adjacent_find_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_binary_search.hpp b/rocprim/include/rocprim/device/detail/config/device_binary_search.hpp index ba9ae94d3..976c9ddec 100644 --- a/rocprim/include/rocprim/device/detail/config/device_binary_search.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_binary_search.hpp @@ -973,6 +973,652 @@ struct default_binary_search_config< : binary_search_config<64, 16> {}; +// Based on value_type = double, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = double, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = double, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = double, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : binary_search_config<128, 1> +{}; + +// Based on value_type = float, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = float, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = float, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = float, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : binary_search_config<128, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : binary_search_config<256, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : binary_search_config<256, 1> +{}; + +// Based on value_type = rocprim::half, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : binary_search_config<128, 2> +{}; + +// Based on value_type = rocprim::half, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 1))>> + : binary_search_config<64, 4> +{}; + +// Based on value_type = int64_t, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : binary_search_config<128, 1> +{}; + +// Based on value_type = int, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : binary_search_config<128, 1> +{}; + +// Based on value_type = short, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<256, 1> +{}; + +// Based on value_type = short, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<256, 1> +{}; + +// Based on value_type = short, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<256, 1> +{}; + +// Based on value_type = short, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 1))>> : binary_search_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : binary_search_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : binary_search_config<128, 1> +{}; + +// Based on value_type = int8_t, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : binary_search_config<64, 2> +{}; + +// Based on value_type = int8_t, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 1))>> + : binary_search_config<64, 8> +{}; + +// Based on value_type = double, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = double, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = double, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = double, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : binary_search_config<256, 1> +{}; + +// Based on value_type = float, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = float, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = float, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = float, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : binary_search_config<128, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : binary_search_config<256, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : binary_search_config<64, 1> +{}; + +// Based on value_type = rocprim::half, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : binary_search_config<64, 2> +{}; + +// Based on value_type = rocprim::half, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : binary_search_config<256, 1> +{}; + +// Based on value_type = int, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<128, 1> +{}; + +// Based on value_type = int, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : binary_search_config<128, 2> +{}; + +// Based on value_type = short, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : binary_search_config<256, 1> +{}; + +// Based on value_type = short, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : binary_search_config<256, 1> +{}; + +// Based on value_type = short, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : binary_search_config<256, 1> +{}; + +// Based on value_type = short, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 1))>> : binary_search_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int64_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : binary_search_config<64, 1> +{}; + +// Based on value_type = int8_t, output_type = int +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : binary_search_config<64, 1> +{}; + +// Based on value_type = int8_t, output_type = short +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : binary_search_config<64, 1> +{}; + +// Based on value_type = int8_t, output_type = int8_t +template +struct default_binary_search_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 1))>> + : binary_search_config<128, 1> +{}; + // Based on value_type = double, output_type = int64_t template struct default_binary_search_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_find_first_of.hpp b/rocprim/include/rocprim/device/detail/config/device_find_first_of.hpp index b908f06fc..b1c8c12b9 100644 --- a/rocprim/include/rocprim/device/detail/config/device_find_first_of.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_find_first_of.hpp @@ -133,6 +133,41 @@ struct default_find_first_of_config< : find_first_of_config<256, 4> {}; +// Based on value_type = int64_t +template +struct default_find_first_of_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<((sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : find_first_of_config<128, 15> +{}; + +// Based on value_type = int +template +struct default_find_first_of_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<((sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : find_first_of_config<256, 12> +{}; + +// Based on value_type = short +template +struct default_find_first_of_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<((sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : find_first_of_config<128, 16> +{}; + +// Based on value_type = int8_t +template +struct default_find_first_of_config(target_arch::gfx1200), + value_type, + std::enable_if_t<((sizeof(value_type) <= 1))>> + : find_first_of_config<256, 16> +{}; + // Based on value_type = int64_t template struct default_find_first_of_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_lower_bound.hpp b/rocprim/include/rocprim/device/detail/config/device_lower_bound.hpp index 2f6f83a9a..e6db00947 100644 --- a/rocprim/include/rocprim/device/detail/config/device_lower_bound.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_lower_bound.hpp @@ -973,6 +973,652 @@ struct default_lower_bound_config< : lower_bound_config<64, 16> {}; +// Based on value_type = double, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : lower_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : lower_bound_config<128, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : lower_bound_config<256, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : lower_bound_config<256, 8> +{}; + +// Based on value_type = rocprim::half, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : lower_bound_config<64, 2> +{}; + +// Based on value_type = rocprim::half, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 1))>> + : lower_bound_config<256, 1> +{}; + +// Based on value_type = int64_t, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : lower_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : lower_bound_config<128, 1> +{}; + +// Based on value_type = short, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 1))>> : lower_bound_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : lower_bound_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : lower_bound_config<128, 4> +{}; + +// Based on value_type = int8_t, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : lower_bound_config<64, 2> +{}; + +// Based on value_type = int8_t, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 1))>> + : lower_bound_config<64, 1> +{}; + +// Based on value_type = double, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : lower_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : lower_bound_config<128, 2> +{}; + +// Based on value_type = rocprim::half, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : lower_bound_config<128, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : lower_bound_config<128, 1> +{}; + +// Based on value_type = rocprim::half, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : lower_bound_config<64, 2> +{}; + +// Based on value_type = rocprim::half, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : lower_bound_config<128, 2> +{}; + +// Based on value_type = int, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : lower_bound_config<128, 2> +{}; + +// Based on value_type = short, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : lower_bound_config<256, 2> +{}; + +// Based on value_type = short, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : lower_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : lower_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 1))>> : lower_bound_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int64_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : lower_bound_config<64, 1> +{}; + +// Based on value_type = int8_t, output_type = int +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : lower_bound_config<128, 4> +{}; + +// Based on value_type = int8_t, output_type = short +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : lower_bound_config<128, 2> +{}; + +// Based on value_type = int8_t, output_type = int8_t +template +struct default_lower_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 1))>> + : lower_bound_config<128, 1> +{}; + // Based on value_type = double, output_type = int64_t template struct default_lower_bound_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_merge.hpp b/rocprim/include/rocprim/device/detail/config/device_merge.hpp index 8e15cbd7f..caada274c 100644 --- a/rocprim/include/rocprim/device/detail/config/device_merge.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_merge.hpp @@ -1129,6 +1129,800 @@ struct default_merge_config< : merge_config<1024, 2> {}; +// Based on key_type = double, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<1024, 2> +{}; + +// Based on key_type = double, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<1024, 4> +{}; + +// Based on key_type = double, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<1024, 4> +{}; + +// Based on key_type = double, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<1024, 4> +{}; + +// Based on key_type = double, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : merge_config<1024, 4> +{}; + +// Based on key_type = float, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<1024, 4> +{}; + +// Based on key_type = float, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<1024, 4> +{}; + +// Based on key_type = float, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<1024, 8> +{}; + +// Based on key_type = float, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = float, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = rocprim::half, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : merge_config<1024, 4> +{}; + +// Based on key_type = rocprim::half, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = rocprim::half, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : merge_config<512, 16> +{}; + +// Based on key_type = rocprim::half, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = rocprim::half, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int64_t, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<1024, 2> +{}; + +// Based on key_type = int64_t, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<1024, 4> +{}; + +// Based on key_type = int64_t, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<1024, 4> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<1024, 4> +{}; + +// Based on key_type = int64_t, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : merge_config<1024, 4> +{}; + +// Based on key_type = int, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<1024, 4> +{}; + +// Based on key_type = int, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<1024, 4> +{}; + +// Based on key_type = int, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<1024, 8> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = int, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = short, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<512, 8> +{}; + +// Based on key_type = short, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<512, 16> +{}; + +// Based on key_type = short, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<512, 16> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = short, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) + && (std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int8_t, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : merge_config<512, 8> +{}; + +// Based on key_type = int8_t, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = int8_t, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int8_t, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = double, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<1024, 2> +{}; + +// Based on key_type = double, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<512, 8> +{}; + +// Based on key_type = double, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<512, 8> +{}; + +// Based on key_type = double, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 8> +{}; + +// Based on key_type = double, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : merge_config<512, 8> +{}; + +// Based on key_type = float, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<1024, 4> +{}; + +// Based on key_type = float, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<1024, 4> +{}; + +// Based on key_type = float, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<512, 16> +{}; + +// Based on key_type = float, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = float, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = rocprim::half, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : merge_config<1024, 4> +{}; + +// Based on key_type = rocprim::half, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = rocprim::half, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : merge_config<512, 16> +{}; + +// Based on key_type = rocprim::half, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = rocprim::half, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int64_t, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<1024, 2> +{}; + +// Based on key_type = int64_t, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<512, 8> +{}; + +// Based on key_type = int64_t, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<512, 8> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 8> +{}; + +// Based on key_type = int64_t, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : merge_config<512, 8> +{}; + +// Based on key_type = int, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<1024, 4> +{}; + +// Based on key_type = int, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<1024, 4> +{}; + +// Based on key_type = int, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<512, 16> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : merge_config<1024, 8> +{}; + +// Based on key_type = short, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_config<512, 8> +{}; + +// Based on key_type = short, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_config<512, 16> +{}; + +// Based on key_type = short, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_config<512, 16> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = short, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) + && (std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int8_t, value_type = int64_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : merge_config<1024, 4> +{}; + +// Based on key_type = int8_t, value_type = int +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : merge_config<512, 8> +{}; + +// Based on key_type = int8_t, value_type = short +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_config<512, 16> +{}; + +// Based on key_type = int8_t, value_type = empty_type +template +struct default_merge_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (std::is_same::value))>> + : merge_config<512, 16> +{}; + // Based on key_type = double, value_type = int64_t template struct default_merge_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_merge_sort_block_sort.hpp b/rocprim/include/rocprim/device/detail/config/device_merge_sort_block_sort.hpp index 59c9287fe..419e362ff 100644 --- a/rocprim/include/rocprim/device/detail/config/device_merge_sort_block_sort.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_merge_sort_block_sort.hpp @@ -1133,6 +1133,954 @@ struct default_merge_sort_block_sort_config< : merge_sort_block_sort_config<256, 32> {}; +// Based on key_type = double, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = double, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<1024, 4> +{}; + +// Based on key_type = double, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<1024, 4> +{}; + +// Based on key_type = double, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = double, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = double, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = float, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = float, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = float, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = float, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = float, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = float, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 16> +{}; + +// Based on key_type = rocprim::half, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 16) && (sizeof(value_type) > 8))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = rocprim::half, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = rocprim::half, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = rocprim::half, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = rocprim::half, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<1024, 8> +{}; + +// Based on key_type = rocprim::half, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 32> +{}; + +// Based on key_type = int64_t, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int64_t, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<1024, 4> +{}; + +// Based on key_type = int64_t, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<1024, 4> +{}; + +// Based on key_type = int64_t, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int64_t, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 16> +{}; + +// Based on key_type = short, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = short, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = short, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<1024, 8> +{}; + +// Based on key_type = short, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<1024, 8> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = short, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 32> +{}; + +// Based on key_type = int8_t, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 16) && (sizeof(value_type) > 8))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int8_t, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int8_t, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int8_t, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int8_t, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 32> +{}; + +// Based on key_type = double, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = double, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<1024, 4> +{}; + +// Based on key_type = double, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = double, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = double, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<1024, 4> +{}; + +// Based on key_type = double, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = float, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = float, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<256, 8> +{}; + +// Based on key_type = float, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = float, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = float, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = float, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 16> +{}; + +// Based on key_type = rocprim::half, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 16) && (sizeof(value_type) > 8))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = rocprim::half, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = rocprim::half, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = rocprim::half, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = rocprim::half, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<256, 4> +{}; + +// Based on key_type = rocprim::half, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (std::is_same::value))>> + : merge_sort_block_sort_config<512, 16> +{}; + +// Based on key_type = int64_t, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int64_t, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<1024, 4> +{}; + +// Based on key_type = int64_t, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int64_t, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int64_t, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 16> +{}; + +// Based on key_type = short, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 16) + && (sizeof(value_type) > 8))>> : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = short, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = short, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : merge_sort_block_sort_config<1024, 8> +{}; + +// Based on key_type = short, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<256, 4> +{}; + +// Based on key_type = short, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 32> +{}; + +// Based on key_type = int8_t, value_type = custom_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 16) && (sizeof(value_type) > 8))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int8_t, value_type = int64_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : merge_sort_block_sort_config<256, 4> +{}; + +// Based on key_type = int8_t, value_type = int +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : merge_sort_block_sort_config<512, 4> +{}; + +// Based on key_type = int8_t, value_type = short +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : merge_sort_block_sort_config<256, 4> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : merge_sort_block_sort_config<512, 8> +{}; + +// Based on key_type = int8_t, value_type = empty_type +template +struct default_merge_sort_block_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (std::is_same::value))>> + : merge_sort_block_sort_config<256, 32> +{}; + // Based on key_type = double, value_type = rocprim::int128_t template struct default_merge_sort_block_sort_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_partition_flag.hpp b/rocprim/include/rocprim/device/detail/config/device_partition_flag.hpp index 3cf0f0d18..62e355c89 100644 --- a/rocprim/include/rocprim/device/detail/config/device_partition_flag.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_partition_flag.hpp @@ -197,6 +197,72 @@ struct default_partition_flag_config< && (sizeof(data_type) <= 1))>> : select_config<128, 28> {}; +// Based on data_type = double +template +struct default_partition_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4))>> : select_config<256, 4> +{}; + +// Based on data_type = float +template +struct default_partition_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2))>> : select_config<384, 6> +{}; + +// Based on data_type = rocprim::half +template +struct default_partition_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2))>> : select_config<256, 12> +{}; + +// Based on data_type = int64_t +template +struct default_partition_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4))>> + : select_config<256, 4> +{}; + +// Based on data_type = int +template +struct default_partition_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2))>> + : select_config<384, 7> +{}; + +// Based on data_type = short +template +struct default_partition_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1))>> + : select_config<256, 12> +{}; + +// Based on data_type = int8_t +template +struct default_partition_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1))>> : select_config<256, 24> +{}; + // Based on data_type = double template struct default_partition_flag_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_partition_predicate.hpp b/rocprim/include/rocprim/device/detail/config/device_partition_predicate.hpp index fdd747029..d69b4ad7c 100644 --- a/rocprim/include/rocprim/device/detail/config/device_partition_predicate.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_partition_predicate.hpp @@ -197,6 +197,72 @@ struct default_partition_predicate_config< && (sizeof(data_type) <= 1))>> : select_config<256, 28> {}; +// Based on data_type = double +template +struct default_partition_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4))>> : select_config<512, 4> +{}; + +// Based on data_type = float +template +struct default_partition_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2))>> : select_config<384, 7> +{}; + +// Based on data_type = rocprim::half +template +struct default_partition_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2))>> : select_config<512, 24> +{}; + +// Based on data_type = int64_t +template +struct default_partition_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4))>> + : select_config<512, 4> +{}; + +// Based on data_type = int +template +struct default_partition_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2))>> + : select_config<384, 7> +{}; + +// Based on data_type = short +template +struct default_partition_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1))>> + : select_config<512, 24> +{}; + +// Based on data_type = int8_t +template +struct default_partition_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1))>> : select_config<512, 28> +{}; + // Based on data_type = double template struct default_partition_predicate_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_partition_three_way.hpp b/rocprim/include/rocprim/device/detail/config/device_partition_three_way.hpp index 6bbe6bd4a..1347cda10 100644 --- a/rocprim/include/rocprim/device/detail/config/device_partition_three_way.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_partition_three_way.hpp @@ -197,6 +197,72 @@ struct default_partition_three_way_config< && (sizeof(data_type) <= 1))>> : select_config<256, 20> {}; +// Based on data_type = double +template +struct default_partition_three_way_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4))>> : select_config<384, 4> +{}; + +// Based on data_type = float +template +struct default_partition_three_way_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2))>> : select_config<512, 10> +{}; + +// Based on data_type = rocprim::half +template +struct default_partition_three_way_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2))>> : select_config<512, 26> +{}; + +// Based on data_type = int64_t +template +struct default_partition_three_way_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4))>> + : select_config<384, 4> +{}; + +// Based on data_type = int +template +struct default_partition_three_way_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2))>> + : select_config<512, 10> +{}; + +// Based on data_type = short +template +struct default_partition_three_way_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1))>> + : select_config<512, 28> +{}; + +// Based on data_type = int8_t +template +struct default_partition_three_way_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1))>> : select_config<256, 28> +{}; + // Based on data_type = double template struct default_partition_three_way_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_partition_two_way_flag.hpp b/rocprim/include/rocprim/device/detail/config/device_partition_two_way_flag.hpp index 87362f2fd..f72284cab 100644 --- a/rocprim/include/rocprim/device/detail/config/device_partition_two_way_flag.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_partition_two_way_flag.hpp @@ -197,6 +197,72 @@ struct default_partition_two_way_flag_config< && (sizeof(data_type) <= 1))>> : select_config<128, 28> {}; +// Based on data_type = double +template +struct default_partition_two_way_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4))>> : select_config<256, 4> +{}; + +// Based on data_type = float +template +struct default_partition_two_way_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2))>> : select_config<384, 7> +{}; + +// Based on data_type = rocprim::half +template +struct default_partition_two_way_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2))>> : select_config<384, 12> +{}; + +// Based on data_type = int64_t +template +struct default_partition_two_way_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4))>> + : select_config<256, 4> +{}; + +// Based on data_type = int +template +struct default_partition_two_way_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2))>> + : select_config<256, 6> +{}; + +// Based on data_type = short +template +struct default_partition_two_way_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1))>> + : select_config<384, 14> +{}; + +// Based on data_type = int8_t +template +struct default_partition_two_way_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1))>> : select_config<256, 24> +{}; + // Based on data_type = double template struct default_partition_two_way_flag_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_partition_two_way_predicate.hpp b/rocprim/include/rocprim/device/detail/config/device_partition_two_way_predicate.hpp index e24efd16d..95a523fcb 100644 --- a/rocprim/include/rocprim/device/detail/config/device_partition_two_way_predicate.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_partition_two_way_predicate.hpp @@ -198,6 +198,72 @@ struct default_partition_two_way_predicate_config< && (sizeof(data_type) <= 1))>> : select_config<256, 28> {}; +// Based on data_type = double +template +struct default_partition_two_way_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4))>> : select_config<256, 4> +{}; + +// Based on data_type = float +template +struct default_partition_two_way_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2))>> : select_config<384, 6> +{}; + +// Based on data_type = rocprim::half +template +struct default_partition_two_way_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2))>> : select_config<384, 18> +{}; + +// Based on data_type = int64_t +template +struct default_partition_two_way_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4))>> + : select_config<512, 4> +{}; + +// Based on data_type = int +template +struct default_partition_two_way_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2))>> + : select_config<512, 8> +{}; + +// Based on data_type = short +template +struct default_partition_two_way_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1))>> + : select_config<384, 18> +{}; + +// Based on data_type = int8_t +template +struct default_partition_two_way_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1))>> : select_config<384, 28> +{}; + // Based on data_type = double template struct default_partition_two_way_predicate_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_radix_sort_block_sort.hpp b/rocprim/include/rocprim/device/detail/config/device_radix_sort_block_sort.hpp index 377aee0d3..e4ece842a 100644 --- a/rocprim/include/rocprim/device/detail/config/device_radix_sort_block_sort.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_radix_sort_block_sort.hpp @@ -1133,6 +1133,403 @@ struct default_radix_sort_block_sort_config< : kernel_config<64, 32> {}; +// Based on key_type = double, value_type = int64_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : kernel_config<64, 28> +{}; + +// Based on key_type = double, value_type = int +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : kernel_config<128, 25> +{}; + +// Based on key_type = double, value_type = short +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : kernel_config<128, 25> +{}; + +// Based on key_type = double, value_type = int8_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : kernel_config<64, 28> +{}; + +// Based on key_type = double, value_type = empty_type +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : kernel_config<128, 32> +{}; + +// Based on key_type = float, value_type = int64_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : kernel_config<128, 25> +{}; + +// Based on key_type = float, value_type = int +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : kernel_config<64, 30> +{}; + +// Based on key_type = float, value_type = short +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : kernel_config<64, 27> +{}; + +// Based on key_type = float, value_type = int8_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : kernel_config<64, 26> +{}; + +// Based on key_type = float, value_type = empty_type +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : kernel_config<128, 32> +{}; + +// Based on key_type = rocprim::half, value_type = int64_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : kernel_config<256, 24> +{}; + +// Based on key_type = rocprim::half, value_type = int +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : kernel_config<128, 29> +{}; + +// Based on key_type = rocprim::half, value_type = short +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : kernel_config<128, 29> +{}; + +// Based on key_type = rocprim::half, value_type = int8_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : kernel_config<128, 27> +{}; + +// Based on key_type = rocprim::half, value_type = empty_type +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (std::is_same::value))>> + : kernel_config<128, 32> +{}; + +// Based on key_type = int64_t, value_type = int64_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : kernel_config<64, 28> +{}; + +// Based on key_type = int64_t, value_type = int +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : kernel_config<128, 30> +{}; + +// Based on key_type = int64_t, value_type = short +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : kernel_config<64, 28> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : kernel_config<64, 28> +{}; + +// Based on key_type = int64_t, value_type = empty_type +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : kernel_config<64, 32> +{}; + +// Based on key_type = int, value_type = int64_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : kernel_config<128, 25> +{}; + +// Based on key_type = int, value_type = int +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : kernel_config<64, 31> +{}; + +// Based on key_type = int, value_type = short +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : kernel_config<64, 30> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : kernel_config<64, 29> +{}; + +// Based on key_type = int, value_type = empty_type +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : kernel_config<128, 32> +{}; + +// Based on key_type = short, value_type = int64_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : kernel_config<256, 24> +{}; + +// Based on key_type = short, value_type = int +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : kernel_config<128, 29> +{}; + +// Based on key_type = short, value_type = short +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : kernel_config<64, 29> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : kernel_config<64, 27> +{}; + +// Based on key_type = short, value_type = empty_type +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) + && (std::is_same::value))>> + : kernel_config<64, 32> +{}; + +// Based on key_type = int8_t, value_type = int64_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : kernel_config<256, 30> +{}; + +// Based on key_type = int8_t, value_type = int +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : kernel_config<256, 29> +{}; + +// Based on key_type = int8_t, value_type = short +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : kernel_config<128, 30> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : kernel_config<64, 27> +{}; + +// Based on key_type = int8_t, value_type = empty_type +template +struct default_radix_sort_block_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (std::is_same::value))>> + : kernel_config<64, 31> +{}; + // Based on key_type = double, value_type = rocprim::int128_t template struct default_radix_sort_block_sort_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_reduce.hpp b/rocprim/include/rocprim/device/detail/config/device_reduce.hpp index d0da09293..48fc7e681 100644 --- a/rocprim/include/rocprim/device/detail/config/device_reduce.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_reduce.hpp @@ -201,6 +201,74 @@ struct default_reduce_config(target_arch::gfx1100), : reduce_config<256, 16, ::rocprim::block_reduce_algorithm::using_warp_reduce> {}; +// Based on key_type = double +template +struct default_reduce_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4))>> + : reduce_config<256, 16, ::rocprim::block_reduce_algorithm::using_warp_reduce> +{}; + +// Based on key_type = float +template +struct default_reduce_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2))>> + : reduce_config<256, 16, ::rocprim::block_reduce_algorithm::using_warp_reduce> +{}; + +// Based on key_type = rocprim::half +template +struct default_reduce_config(target_arch::gfx1200), + key_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(key_type) <= 2))>> + : reduce_config<256, 16, ::rocprim::block_reduce_algorithm::using_warp_reduce> +{}; + +// Based on key_type = int64_t +template +struct default_reduce_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4))>> + : reduce_config<256, 16, ::rocprim::block_reduce_algorithm::using_warp_reduce> +{}; + +// Based on key_type = int +template +struct default_reduce_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2))>> + : reduce_config<256, 16, ::rocprim::block_reduce_algorithm::using_warp_reduce> +{}; + +// Based on key_type = short +template +struct default_reduce_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1))>> + : reduce_config<256, 16, ::rocprim::block_reduce_algorithm::using_warp_reduce> +{}; + +// Based on key_type = int8_t +template +struct default_reduce_config(target_arch::gfx1200), + key_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(key_type) <= 1))>> + : reduce_config<256, 16, ::rocprim::block_reduce_algorithm::using_warp_reduce> +{}; + // Based on key_type = double template struct default_reduce_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_reduce_by_key.hpp b/rocprim/include/rocprim/device/detail/config/device_reduce_by_key.hpp index c63e26316..92edbc6e6 100644 --- a/rocprim/include/rocprim/device/detail/config/device_reduce_by_key.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_reduce_by_key.hpp @@ -946,6 +946,451 @@ struct default_reduce_by_key_config< block_scan_algorithm::using_warp_scan> {}; +// Based on key_type = int64_t, value_type = double +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 8) && (sizeof(key_type) > 4) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : reduce_by_key_config<512, + 6, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t, value_type = float +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 8) && (sizeof(key_type) > 4) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : reduce_by_key_config<512, + 4, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t, value_type = rocprim::half +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 8) && (sizeof(key_type) > 4) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2))>> + : reduce_by_key_config<512, + 6, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t, value_type = int64_t +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 8) && (sizeof(key_type) > 4) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : reduce_by_key_config<512, + 11, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t, value_type = int +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 8) && (sizeof(key_type) > 4) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : reduce_by_key_config<512, + 4, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t, value_type = short +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 8) && (sizeof(key_type) > 4) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : reduce_by_key_config<512, + 6, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 8) && (sizeof(key_type) > 4) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1))>> + : reduce_by_key_config<256, + 10, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = double +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 4) && (sizeof(key_type) > 2) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : reduce_by_key_config<256, + 8, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = float +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 4) && (sizeof(key_type) > 2) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : reduce_by_key_config<512, + 8, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = rocprim::half +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 4) && (sizeof(key_type) > 2) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2))>> + : reduce_by_key_config<512, + 12, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = int64_t +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 4) && (sizeof(key_type) > 2) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : reduce_by_key_config<512, + 8, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = int +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 4) && (sizeof(key_type) > 2) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : reduce_by_key_config<512, + 8, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = short +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 4) && (sizeof(key_type) > 2) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : reduce_by_key_config<512, + 12, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 4) && (sizeof(key_type) > 2) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1))>> + : reduce_by_key_config<512, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short, value_type = double +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 2) && (sizeof(key_type) > 1) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : reduce_by_key_config<256, + 9, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short, value_type = float +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 2) && (sizeof(key_type) > 1) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : reduce_by_key_config<512, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short, value_type = rocprim::half +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 2) && (sizeof(key_type) > 1) + && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2))>> + : reduce_by_key_config<384, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short, value_type = int64_t +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 2) && (sizeof(key_type) > 1) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : reduce_by_key_config<384, + 8, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short, value_type = int +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 2) && (sizeof(key_type) > 1) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : reduce_by_key_config<512, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short, value_type = short +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 2) && (sizeof(key_type) > 1) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : reduce_by_key_config<384, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 2) && (sizeof(key_type) > 1) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1))>> + : reduce_by_key_config<384, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = double +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 1) && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : reduce_by_key_config<192, + 12, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = float +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 1) && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = rocprim::half +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 1) && bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = int64_t +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 1) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : reduce_by_key_config<384, + 8, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = int +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 1) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = short +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 1) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_reduce_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<((sizeof(key_type) <= 1) + && !bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_transpose, + block_scan_algorithm::using_warp_scan> +{}; + // Based on key_type = int64_t, value_type = double template struct default_reduce_by_key_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_run_length_encode.hpp b/rocprim/include/rocprim/device/detail/config/device_run_length_encode.hpp index 57c5efc03..ef625ee73 100644 --- a/rocprim/include/rocprim/device/detail/config/device_run_length_encode.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_run_length_encode.hpp @@ -283,6 +283,216 @@ struct default_trivial_runs_config< block_scan_algorithm::using_warp_scan> {}; +// Based on key_type = double +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4))>> + : reduce_by_key_config<512, + 7, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = float +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2))>> + : reduce_by_key_config<256, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = rocprim::half +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(key_type) <= 2))>> + : reduce_by_key_config<384, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4))>> + : reduce_by_key_config<512, + 7, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2))>> + : reduce_by_key_config<256, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1))>> + : reduce_by_key_config<384, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(key_type) <= 1))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = double +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4))>> + : reduce_by_key_config<512, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = float +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = rocprim::half +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(key_type) <= 2))>> + : reduce_by_key_config<512, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4))>> + : reduce_by_key_config<512, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1))>> + : reduce_by_key_config<512, + 14, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t +template +struct default_trivial_runs_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(key_type) <= 1))>> + : reduce_by_key_config<384, + 15, + block_load_method::block_load_transpose, + block_load_method::block_load_direct, + block_scan_algorithm::using_warp_scan> +{}; + // Based on key_type = double template struct default_trivial_runs_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_run_length_encode_non_trivial.hpp b/rocprim/include/rocprim/device/detail/config/device_run_length_encode_non_trivial.hpp index 9b304330a..ff7fde76f 100644 --- a/rocprim/include/rocprim/device/detail/config/device_run_length_encode_non_trivial.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_run_length_encode_non_trivial.hpp @@ -253,6 +253,97 @@ struct default_non_trivial_runs_config< ::rocprim::block_scan_algorithm::using_warp_scan> {}; +// Based on key_type = double +template +struct default_non_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4))>> + : non_trivial_runs_config<256, + 8, + ::rocprim::block_load_method::block_load_vectorize, + ::rocprim::block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = float +template +struct default_non_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2))>> + : non_trivial_runs_config<256, + 8, + ::rocprim::block_load_method::block_load_vectorize, + ::rocprim::block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = rocprim::half +template +struct default_non_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(key_type) <= 2))>> + : non_trivial_runs_config<128, + 32, + ::rocprim::block_load_method::block_load_vectorize, + ::rocprim::block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int64_t +template +struct default_non_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4))>> + : non_trivial_runs_config<256, + 8, + ::rocprim::block_load_method::block_load_vectorize, + ::rocprim::block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int +template +struct default_non_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2))>> + : non_trivial_runs_config<256, + 8, + ::rocprim::block_load_method::block_load_vectorize, + ::rocprim::block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short +template +struct default_non_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1))>> + : non_trivial_runs_config<128, + 16, + ::rocprim::block_load_method::block_load_vectorize, + ::rocprim::block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t +template +struct default_non_trivial_runs_config< + static_cast(target_arch::gfx1200), + key_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(key_type) <= 1))>> + : non_trivial_runs_config<128, + 32, + ::rocprim::block_load_method::block_load_vectorize, + ::rocprim::block_scan_algorithm::using_warp_scan> +{}; + // Based on key_type = double template struct default_non_trivial_runs_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_scan.hpp b/rocprim/include/rocprim/device/detail/config/device_scan.hpp index 2bf468c29..3d3990537 100644 --- a/rocprim/include/rocprim/device/detail/config/device_scan.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_scan.hpp @@ -265,6 +265,198 @@ struct default_scan_config(target_arch::gfx1100), block_scan_algorithm::using_warp_scan> {}; +// Based on value_type = double +template +struct default_scan_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : scan_config<256, + 24, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on value_type = float +template +struct default_scan_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : scan_config<256, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on value_type = rocprim::half +template +struct default_scan_config(target_arch::gfx1200), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2))>> + : scan_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on value_type = int64_t +template +struct default_scan_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : scan_config<256, + 4, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on value_type = int +template +struct default_scan_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : scan_config<256, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on value_type = short +template +struct default_scan_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : scan_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on value_type = int8_t +template +struct default_scan_config(target_arch::gfx1200), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1))>> + : scan_config<256, + 24, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on value_type = double +template +struct default_scan_config< + static_cast(target_arch::gfx1201), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : scan_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on value_type = float +template +struct default_scan_config< + static_cast(target_arch::gfx1201), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : scan_config<256, + 6, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on value_type = rocprim::half +template +struct default_scan_config(target_arch::gfx1201), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2))>> + : scan_config<256, + 12, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on value_type = int64_t +template +struct default_scan_config< + static_cast(target_arch::gfx1201), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : scan_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on value_type = int +template +struct default_scan_config< + static_cast(target_arch::gfx1201), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : scan_config<256, + 6, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on value_type = short +template +struct default_scan_config< + static_cast(target_arch::gfx1201), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : scan_config<256, + 12, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on value_type = int8_t +template +struct default_scan_config(target_arch::gfx1201), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1))>> + : scan_config<256, + 24, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + // Based on value_type = double template struct default_scan_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_scan_by_key.hpp b/rocprim/include/rocprim/device/detail/config/device_scan_by_key.hpp index 1a380f8b7..479bac10e 100644 --- a/rocprim/include/rocprim/device/detail/config/device_scan_by_key.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_scan_by_key.hpp @@ -1293,6 +1293,441 @@ struct default_scan_by_key_config< block_scan_algorithm::using_warp_scan> {}; +// Based on key_type = double, value_type = int64_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> + : scan_by_key_config<256, + 4, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = double, value_type = int +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> + : scan_by_key_config<256, + 20, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = double, value_type = short +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> + : scan_by_key_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = double, value_type = int8_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>> + : scan_by_key_config<128, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = float, value_type = int64_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> + : scan_by_key_config<256, + 4, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = float, value_type = int +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> + : scan_by_key_config<256, + 5, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = float, value_type = short +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> + : scan_by_key_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = float, value_type = int8_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>> + : scan_by_key_config<256, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = rocprim::half, value_type = int64_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : scan_by_key_config<256, + 6, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = rocprim::half, value_type = int +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : scan_by_key_config<256, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = rocprim::half, value_type = short +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : scan_by_key_config<256, + 10, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = rocprim::half, value_type = int8_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 1))>> + : scan_by_key_config<256, + 20, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = int64_t, value_type = int64_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> + : scan_by_key_config<256, + 4, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = int64_t, value_type = int +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> + : scan_by_key_config<256, + 20, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = int64_t, value_type = short +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> + : scan_by_key_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>> + : scan_by_key_config<128, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = int64_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> + : scan_by_key_config<256, + 4, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = int +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> + : scan_by_key_config<256, + 5, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = int, value_type = short +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> + : scan_by_key_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>> + : scan_by_key_config<256, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = short, value_type = int64_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> + : scan_by_key_config<256, + 6, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = short, value_type = int +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> + : scan_by_key_config<256, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = short, value_type = short +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> + : scan_by_key_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1))>> + : scan_by_key_config<256, + 20, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + +// Based on key_type = int8_t, value_type = int64_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : scan_by_key_config<256, + 12, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = int +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : scan_by_key_config<256, + 8, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = short +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : scan_by_key_config<256, + 16, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::using_warp_scan> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_scan_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 1))>> + : scan_by_key_config<256, + 24, + ::rocprim::block_load_method::block_load_transpose, + ::rocprim::block_store_method::block_store_transpose, + block_scan_algorithm::reduce_then_scan> +{}; + // Based on key_type = double, value_type = rocprim::int128_t template struct default_scan_by_key_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_segmented_radix_sort.hpp b/rocprim/include/rocprim/device/detail/config/device_segmented_radix_sort.hpp index b1ae4e50f..1003b35e7 100644 --- a/rocprim/include/rocprim/device/detail/config/device_segmented_radix_sort.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_segmented_radix_sort.hpp @@ -1849,6 +1849,270 @@ struct default_segmented_radix_sort_config< 1> {}; +// Based on key_type = double, value_type = empty_type +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 0, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = float, value_type = empty_type +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 0, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = rocprim::half, value_type = empty_type +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 0, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = int64_t, value_type = empty_type +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) + && (std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 0, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = int, value_type = empty_type +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) + && (std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 0, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = short, value_type = empty_type +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) + && (std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 0, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = int8_t, value_type = empty_type +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 0, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = double, value_type = int8_t +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 8, + kernel_config<256, 4>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = float, value_type = int8_t +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 8, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = rocprim::half, value_type = int8_t +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 8, + kernel_config<256, 4>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 8, + kernel_config<256, 4>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 8, + kernel_config<256, 4>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 8, + kernel_config<256, 4>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_segmented_radix_sort_config< + static_cast(target_arch::gfx1201), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 1) + && (!std::is_same::value))>> + : segmented_radix_sort_config< + 8, + 8, + kernel_config<256, 8>, + typename std::conditional<1, + WarpSortConfig<8, 4, 256, 64, 16, 8, 256>, + DisabledWarpSortConfig>::type, + 1> +{}; + // Based on key_type = double, value_type = rocprim::int128_t template struct default_segmented_radix_sort_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_select_flag.hpp b/rocprim/include/rocprim/device/detail/config/device_select_flag.hpp index f17a92893..2c4e035e0 100644 --- a/rocprim/include/rocprim/device/detail/config/device_select_flag.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_select_flag.hpp @@ -197,6 +197,72 @@ struct default_select_flag_config< && (sizeof(data_type) <= 1))>> : select_config<256, 28> {}; +// Based on data_type = double +template +struct default_select_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4))>> : select_config<256, 5> +{}; + +// Based on data_type = float +template +struct default_select_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2))>> : select_config<384, 7> +{}; + +// Based on data_type = rocprim::half +template +struct default_select_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2))>> : select_config<256, 12> +{}; + +// Based on data_type = int64_t +template +struct default_select_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4))>> + : select_config<256, 4> +{}; + +// Based on data_type = int +template +struct default_select_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2))>> + : select_config<192, 8> +{}; + +// Based on data_type = short +template +struct default_select_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1))>> + : select_config<256, 12> +{}; + +// Based on data_type = int8_t +template +struct default_select_flag_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1))>> : select_config<384, 28> +{}; + // Based on data_type = double template struct default_select_flag_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_select_predicate.hpp b/rocprim/include/rocprim/device/detail/config/device_select_predicate.hpp index 2d7c464cd..dd69d9cb2 100644 --- a/rocprim/include/rocprim/device/detail/config/device_select_predicate.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_select_predicate.hpp @@ -197,6 +197,72 @@ struct default_select_predicate_config< && (sizeof(data_type) <= 1))>> : select_config<256, 28> {}; +// Based on data_type = double +template +struct default_select_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4))>> : select_config<256, 4> +{}; + +// Based on data_type = float +template +struct default_select_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2))>> : select_config<384, 7> +{}; + +// Based on data_type = rocprim::half +template +struct default_select_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2))>> : select_config<384, 18> +{}; + +// Based on data_type = int64_t +template +struct default_select_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4))>> + : select_config<256, 4> +{}; + +// Based on data_type = int +template +struct default_select_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2))>> + : select_config<384, 7> +{}; + +// Based on data_type = short +template +struct default_select_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1))>> + : select_config<384, 18> +{}; + +// Based on data_type = int8_t +template +struct default_select_predicate_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1))>> : select_config<384, 28> +{}; + // Based on data_type = double template struct default_select_predicate_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_select_predicated_flag.hpp b/rocprim/include/rocprim/device/detail/config/device_select_predicated_flag.hpp index d61a3433e..699ff1fb0 100644 --- a/rocprim/include/rocprim/device/detail/config/device_select_predicated_flag.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_select_predicated_flag.hpp @@ -955,6 +955,322 @@ struct default_select_predicated_flag_config< : select_config<512, 20> {}; +// Based on data_type = double, flag_type = int64_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4) && (sizeof(flag_type) <= 8) + && (sizeof(flag_type) > 4))>> : select_config<192, 4> +{}; + +// Based on data_type = double, flag_type = int +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4) && (sizeof(flag_type) <= 4) + && (sizeof(flag_type) > 2))>> : select_config<256, 4> +{}; + +// Based on data_type = double, flag_type = short +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4) && (sizeof(flag_type) <= 2) + && (sizeof(flag_type) > 1))>> : select_config<256, 4> +{}; + +// Based on data_type = double, flag_type = int8_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4) && (sizeof(flag_type) <= 1))>> + : select_config<256, 4> +{}; + +// Based on data_type = float, flag_type = int64_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2) && (sizeof(flag_type) <= 8) + && (sizeof(flag_type) > 4))>> : select_config<384, 4> +{}; + +// Based on data_type = float, flag_type = int +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2) && (sizeof(flag_type) <= 4) + && (sizeof(flag_type) > 2))>> : select_config<384, 4> +{}; + +// Based on data_type = float, flag_type = short +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2) && (sizeof(flag_type) <= 2) + && (sizeof(flag_type) > 1))>> : select_config<384, 6> +{}; + +// Based on data_type = float, flag_type = int8_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2) && (sizeof(flag_type) <= 1))>> + : select_config<384, 7> +{}; + +// Based on data_type = rocprim::half, flag_type = int64_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 2) + && (sizeof(flag_type) <= 8) && (sizeof(flag_type) > 4))>> + : select_config<256, 4> +{}; + +// Based on data_type = rocprim::half, flag_type = int +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 2) + && (sizeof(flag_type) <= 4) && (sizeof(flag_type) > 2))>> + : select_config<384, 8> +{}; + +// Based on data_type = rocprim::half, flag_type = short +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 2) + && (sizeof(flag_type) <= 2) && (sizeof(flag_type) > 1))>> + : select_config<384, 12> +{}; + +// Based on data_type = rocprim::half, flag_type = int8_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 2) + && (sizeof(flag_type) <= 1))>> : select_config<384, 18> +{}; + +// Based on data_type = int64_t, flag_type = int64_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4) + && (sizeof(flag_type) <= 8) && (sizeof(flag_type) > 4))>> + : select_config<192, 4> +{}; + +// Based on data_type = int64_t, flag_type = int +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4) + && (sizeof(flag_type) <= 4) && (sizeof(flag_type) > 2))>> + : select_config<256, 4> +{}; + +// Based on data_type = int64_t, flag_type = short +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4) + && (sizeof(flag_type) <= 2) && (sizeof(flag_type) > 1))>> + : select_config<256, 4> +{}; + +// Based on data_type = int64_t, flag_type = int8_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4) + && (sizeof(flag_type) <= 1))>> : select_config<256, 4> +{}; + +// Based on data_type = int, flag_type = int64_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2) + && (sizeof(flag_type) <= 8) && (sizeof(flag_type) > 4))>> + : select_config<384, 4> +{}; + +// Based on data_type = int, flag_type = int +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2) + && (sizeof(flag_type) <= 4) && (sizeof(flag_type) > 2))>> + : select_config<384, 4> +{}; + +// Based on data_type = int, flag_type = short +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2) + && (sizeof(flag_type) <= 2) && (sizeof(flag_type) > 1))>> + : select_config<384, 6> +{}; + +// Based on data_type = int, flag_type = int8_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2) + && (sizeof(flag_type) <= 1))>> : select_config<384, 7> +{}; + +// Based on data_type = short, flag_type = int64_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1) + && (sizeof(flag_type) <= 8) && (sizeof(flag_type) > 4))>> + : select_config<256, 4> +{}; + +// Based on data_type = short, flag_type = int +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1) + && (sizeof(flag_type) <= 4) && (sizeof(flag_type) > 2))>> + : select_config<384, 8> +{}; + +// Based on data_type = short, flag_type = short +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1) + && (sizeof(flag_type) <= 2) && (sizeof(flag_type) > 1))>> + : select_config<384, 12> +{}; + +// Based on data_type = short, flag_type = int8_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1) + && (sizeof(flag_type) <= 1))>> : select_config<384, 18> +{}; + +// Based on data_type = int8_t, flag_type = int64_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1) && (sizeof(flag_type) <= 8) + && (sizeof(flag_type) > 4))>> : select_config<256, 8> +{}; + +// Based on data_type = int8_t, flag_type = int +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1) && (sizeof(flag_type) <= 4) + && (sizeof(flag_type) > 2))>> : select_config<512, 16> +{}; + +// Based on data_type = int8_t, flag_type = short +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1) && (sizeof(flag_type) <= 2) + && (sizeof(flag_type) > 1))>> : select_config<512, 16> +{}; + +// Based on data_type = int8_t, flag_type = int8_t +template +struct default_select_predicated_flag_config< + static_cast(target_arch::gfx1200), + data_type, + flag_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1) && (sizeof(flag_type) <= 1))>> + : select_config<256, 24> +{}; + // Based on data_type = double, flag_type = rocprim::int128_t template struct default_select_predicated_flag_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_select_unique.hpp b/rocprim/include/rocprim/device/detail/config/device_select_unique.hpp index 6653818e7..670ae5da6 100644 --- a/rocprim/include/rocprim/device/detail/config/device_select_unique.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_select_unique.hpp @@ -197,6 +197,72 @@ struct default_select_unique_config< && (sizeof(data_type) <= 1))>> : select_config<384, 28> {}; +// Based on data_type = double +template +struct default_select_unique_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 8) + && (sizeof(data_type) > 4))>> : select_config<384, 8> +{}; + +// Based on data_type = float +template +struct default_select_unique_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(data_type) <= 4) + && (sizeof(data_type) > 2))>> : select_config<384, 9> +{}; + +// Based on data_type = rocprim::half +template +struct default_select_unique_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2))>> : select_config<256, 21> +{}; + +// Based on data_type = int64_t +template +struct default_select_unique_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 8) && (sizeof(data_type) > 4))>> + : select_config<384, 7> +{}; + +// Based on data_type = int +template +struct default_select_unique_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 4) && (sizeof(data_type) > 2))>> + : select_config<384, 9> +{}; + +// Based on data_type = short +template +struct default_select_unique_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 2) && (sizeof(data_type) > 1))>> + : select_config<256, 21> +{}; + +// Based on data_type = int8_t +template +struct default_select_unique_config< + static_cast(target_arch::gfx1200), + data_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(data_type) <= 1))>> : select_config<256, 20> +{}; + // Based on data_type = double template struct default_select_unique_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_select_unique_by_key.hpp b/rocprim/include/rocprim/device/detail/config/device_select_unique_by_key.hpp index 36c641f00..ce67e878e 100644 --- a/rocprim/include/rocprim/device/detail/config/device_select_unique_by_key.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_select_unique_by_key.hpp @@ -921,6 +921,312 @@ struct default_select_unique_by_key_config< && (sizeof(value_type) <= 1))>> : select_config<384, 20> {}; +// Based on key_type = double, value_type = int64_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : select_config<384, 8> +{}; + +// Based on key_type = double, value_type = int +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : select_config<384, 6> +{}; + +// Based on key_type = double, value_type = short +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : select_config<384, 6> +{}; + +// Based on key_type = double, value_type = int8_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>> + : select_config<384, 6> +{}; + +// Based on key_type = float, value_type = int64_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : select_config<384, 4> +{}; + +// Based on key_type = float, value_type = int +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : select_config<384, 7> +{}; + +// Based on key_type = float, value_type = short +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : select_config<384, 9> +{}; + +// Based on key_type = float, value_type = int8_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>> + : select_config<384, 12> +{}; + +// Based on key_type = rocprim::half, value_type = int64_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : select_config<384, 8> +{}; + +// Based on key_type = rocprim::half, value_type = int +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : select_config<512, 16> +{}; + +// Based on key_type = rocprim::half, value_type = short +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : select_config<512, 14> +{}; + +// Based on key_type = rocprim::half, value_type = int8_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(value_type) <= 1))>> : select_config<512, 14> +{}; + +// Based on key_type = int64_t, value_type = int64_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : select_config<384, 8> +{}; + +// Based on key_type = int64_t, value_type = int +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : select_config<384, 6> +{}; + +// Based on key_type = int64_t, value_type = short +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : select_config<384, 6> +{}; + +// Based on key_type = int64_t, value_type = int8_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 8) + && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1))>> + : select_config<384, 8> +{}; + +// Based on key_type = int, value_type = int64_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : select_config<384, 4> +{}; + +// Based on key_type = int, value_type = int +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : select_config<384, 7> +{}; + +// Based on key_type = int, value_type = short +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : select_config<384, 9> +{}; + +// Based on key_type = int, value_type = int8_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 4) + && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1))>> + : select_config<384, 12> +{}; + +// Based on key_type = short, value_type = int64_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8) + && (sizeof(value_type) > 4))>> : select_config<384, 6> +{}; + +// Based on key_type = short, value_type = int +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4) + && (sizeof(value_type) > 2))>> : select_config<512, 16> +{}; + +// Based on key_type = short, value_type = short +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2) + && (sizeof(value_type) > 1))>> : select_config<512, 24> +{}; + +// Based on key_type = short, value_type = int8_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 2) + && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1))>> + : select_config<512, 14> +{}; + +// Based on key_type = int8_t, value_type = int64_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : select_config<256, 8> +{}; + +// Based on key_type = int8_t, value_type = int +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : select_config<512, 16> +{}; + +// Based on key_type = int8_t, value_type = short +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : select_config<512, 32> +{}; + +// Based on key_type = int8_t, value_type = int8_t +template +struct default_select_unique_by_key_config< + static_cast(target_arch::gfx1200), + key_type, + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) && (sizeof(key_type) <= 1) + && (sizeof(value_type) <= 1))>> : select_config<384, 16> +{}; + // Based on key_type = double, value_type = rocprim::int128_t template struct default_select_unique_by_key_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_transform.hpp b/rocprim/include/rocprim/device/detail/config/device_transform.hpp index 7a60e3116..8490901a4 100644 --- a/rocprim/include/rocprim/device/detail/config/device_transform.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_transform.hpp @@ -201,6 +201,74 @@ struct default_transform_config< && (sizeof(value_type) <= 1))>> : transform_config<1024, 4> {}; +// Based on value_type = double +template +struct default_transform_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : transform_config<512, 2> +{}; + +// Based on value_type = float +template +struct default_transform_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : transform_config<512, 1> +{}; + +// Based on value_type = rocprim::half +template +struct default_transform_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2))>> : transform_config<512, 2> +{}; + +// Based on value_type = int64_t +template +struct default_transform_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>> + : transform_config<512, 2> +{}; + +// Based on value_type = int +template +struct default_transform_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>> + : transform_config<512, 1> +{}; + +// Based on value_type = short +template +struct default_transform_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>> + : transform_config<512, 2> +{}; + +// Based on value_type = int8_t +template +struct default_transform_config< + static_cast(target_arch::gfx1200), + value_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1))>> : transform_config<256, 4> +{}; + // Based on value_type = double template struct default_transform_config< diff --git a/rocprim/include/rocprim/device/detail/config/device_upper_bound.hpp b/rocprim/include/rocprim/device/detail/config/device_upper_bound.hpp index ff1ad05d4..743a656d7 100644 --- a/rocprim/include/rocprim/device/detail/config/device_upper_bound.hpp +++ b/rocprim/include/rocprim/device/detail/config/device_upper_bound.hpp @@ -973,6 +973,652 @@ struct default_upper_bound_config< : upper_bound_config<128, 16> {}; +// Based on value_type = double, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : upper_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : upper_bound_config<128, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : upper_bound_config<256, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : upper_bound_config<256, 8> +{}; + +// Based on value_type = rocprim::half, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : upper_bound_config<64, 2> +{}; + +// Based on value_type = rocprim::half, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : upper_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : upper_bound_config<128, 1> +{}; + +// Based on value_type = short, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 1))>> : upper_bound_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : upper_bound_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : upper_bound_config<128, 4> +{}; + +// Based on value_type = int8_t, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : upper_bound_config<128, 2> +{}; + +// Based on value_type = int8_t, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1200), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 1))>> + : upper_bound_config<256, 1> +{}; + +// Based on value_type = double, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = double, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : upper_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = float, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : upper_bound_config<128, 2> +{}; + +// Based on value_type = rocprim::half, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : upper_bound_config<128, 1> +{}; + +// Based on value_type = rocprim::half, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : upper_bound_config<64, 1> +{}; + +// Based on value_type = rocprim::half, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : upper_bound_config<128, 2> +{}; + +// Based on value_type = rocprim::half, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(output_type) <= 1))>> + : upper_bound_config<64, 1> +{}; + +// Based on value_type = int64_t, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int64_t, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4) + && (sizeof(output_type) <= 1))>> : upper_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<128, 1> +{}; + +// Based on value_type = int, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2) + && (sizeof(output_type) <= 1))>> : upper_bound_config<128, 2> +{}; + +// Based on value_type = short, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 8) && (sizeof(output_type) > 4))>> + : upper_bound_config<256, 2> +{}; + +// Based on value_type = short, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 4) && (sizeof(output_type) > 2))>> + : upper_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 2) && (sizeof(output_type) > 1))>> + : upper_bound_config<256, 1> +{}; + +// Based on value_type = short, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1) + && (sizeof(output_type) <= 1))>> : upper_bound_config<256, 1> +{}; + +// Based on value_type = int8_t, output_type = int64_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 8) + && (sizeof(output_type) > 4))>> : upper_bound_config<128, 1> +{}; + +// Based on value_type = int8_t, output_type = int +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 4) + && (sizeof(output_type) > 2))>> : upper_bound_config<128, 4> +{}; + +// Based on value_type = int8_t, output_type = short +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 2) + && (sizeof(output_type) > 1))>> : upper_bound_config<64, 2> +{}; + +// Based on value_type = int8_t, output_type = int8_t +template +struct default_upper_bound_config< + static_cast(target_arch::gfx1201), + value_type, + output_type, + std::enable_if_t<(!bool(rocprim::is_floating_point::value) + && (sizeof(value_type) <= 1) && (sizeof(output_type) <= 1))>> + : upper_bound_config<64, 1> +{}; + // Based on value_type = double, output_type = int64_t template struct default_upper_bound_config<