// Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in
// all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
// THE SOFTWARE.

#ifndef ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_RADIX_SORT_BLOCK_SORT_HPP_
#define ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_RADIX_SORT_BLOCK_SORT_HPP_

#include "../../../type_traits.hpp"
#include "../device_config_helper.hpp"
#include <type_traits>

/* DO NOT EDIT THIS FILE
 * This file is automatically generated by `/scripts/autotune/create_optimization.py`.
 * so most likely you want to edit rocprim/device/device_(algo)_config.hpp
 */

/// \addtogroup primitivesmodule_deviceconfigs
/// @{

BEGIN_ROCPRIM_NAMESPACE

namespace detail
{

template<unsigned int arch,
         class key_type,
         class value_type = rocprim::empty_type,
         class enable     = void>
struct default_radix_sort_block_sort_config
    : radix_sort_block_sort_config_base<key_type, value_type>::type
{};

// Based on key_type = double, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 31>
{};

// Based on key_type = double, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 31>
{};

// Based on key_type = double, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 31>
{};

// Based on key_type = double, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 31>
{};

// Based on key_type = double, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 31>
{};

// Based on key_type = float, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 31>
{};

// Based on key_type = float, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 32>
{};

// Based on key_type = float, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 32>
{};

// Based on key_type = float, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

// Based on key_type = float, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<128, 32>
{};

// Based on key_type = rocprim::half, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<512, 11>
{};

// Based on key_type = rocprim::half, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<1024, 14>
{};

// Based on key_type = rocprim::half, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<512, 8>
{};

// Based on key_type = rocprim::half, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 18>
{};

// Based on key_type = rocprim::half, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 14>
{};

// Based on key_type = int64_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 31>
{};

// Based on key_type = int64_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 31>
{};

// Based on key_type = int64_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 31>
{};

// Based on key_type = int64_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 31>
{};

// Based on key_type = int64_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 31>
{};

// Based on key_type = int, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 31>
{};

// Based on key_type = int, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 32>
{};

// Based on key_type = int, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 32>
{};

// Based on key_type = int, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

// Based on key_type = int, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<128, 32>
{};

// Based on key_type = short, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 16>
{};

// Based on key_type = short, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 32>
{};

// Based on key_type = short, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 32>
{};

// Based on key_type = short, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

// Based on key_type = short, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 32>
{};

// Based on key_type = int8_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<1024, 7>
{};

// Based on key_type = int8_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = int8_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<256, 32>
{};

// Based on key_type = int8_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 28>
{};

// Based on key_type = int8_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx1030),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 32>
{};

// Based on key_type = double, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<128, 13>
{};

// Based on key_type = double, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 13>
{};

// Based on key_type = double, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 14>
{};

// Based on key_type = double, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = double, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = float, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = float, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<128, 17>
{};

// Based on key_type = float, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 15>
{};

// Based on key_type = float, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 17>
{};

// Based on key_type = float, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 21>
{};

// Based on key_type = rocprim::half, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<512, 5>
{};

// Based on key_type = rocprim::half, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<256, 9>
{};

// Based on key_type = rocprim::half, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<256, 7>
{};

// Based on key_type = rocprim::half, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 10>
{};

// Based on key_type = rocprim::half, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 12>
{};

// Based on key_type = int64_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<64, 29>
{};

// Based on key_type = int64_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 13>
{};

// Based on key_type = int64_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 14>
{};

// Based on key_type = int64_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int64_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = int, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<128, 17>
{};

// Based on key_type = int, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 16>
{};

// Based on key_type = int, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 18>
{};

// Based on key_type = int, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 21>
{};

// Based on key_type = short, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<512, 5>
{};

// Based on key_type = short, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 9>
{};

// Based on key_type = short, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 7>
{};

// Based on key_type = short, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 10>
{};

// Based on key_type = short, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

// Based on key_type = int8_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<1024, 6>
{};

// Based on key_type = int8_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<512, 7>
{};

// Based on key_type = int8_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<512, 7>
{};

// Based on key_type = int8_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 20>
{};

// Based on key_type = int8_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx900),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 32>
{};

// Based on key_type = double, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = double, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 13>
{};

// Based on key_type = double, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 14>
{};

// Based on key_type = double, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<128, 15>
{};

// Based on key_type = double, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = float, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = float, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<512, 31>
{};

// Based on key_type = float, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<1024, 15>
{};

// Based on key_type = float, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = float, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 31>
{};

// Based on key_type = rocprim::half, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<256, 7>
{};

// Based on key_type = rocprim::half, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<256, 9>
{};

// Based on key_type = rocprim::half, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<256, 7>
{};

// Based on key_type = rocprim::half, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 10>
{};

// Based on key_type = rocprim::half, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 12>
{};

// Based on key_type = int64_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = int64_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 13>
{};

// Based on key_type = int64_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 14>
{};

// Based on key_type = int64_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int64_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = int, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<512, 31>
{};

// Based on key_type = int, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<1024, 15>
{};

// Based on key_type = int, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = int, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = short, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<512, 13>
{};

// Based on key_type = short, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<1024, 13>
{};

// Based on key_type = short, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<64, 32>
{};

// Based on key_type = short, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 32>
{};

// Based on key_type = short, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

// Based on key_type = int8_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<512, 14>
{};

// Based on key_type = int8_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<512, 24>
{};

// Based on key_type = int8_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<256, 7>
{};

// Based on key_type = int8_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 20>
{};

// Based on key_type = int8_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx908),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 25>
{};

// Based on key_type = double, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = double, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 13>
{};

// Based on key_type = double, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 14>
{};

// Based on key_type = double, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<128, 15>
{};

// Based on key_type = double, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = float, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = float, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<512, 31>
{};

// Based on key_type = float, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<1024, 15>
{};

// Based on key_type = float, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = float, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 31>
{};

// Based on key_type = rocprim::half, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<256, 7>
{};

// Based on key_type = rocprim::half, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<256, 9>
{};

// Based on key_type = rocprim::half, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<256, 7>
{};

// Based on key_type = rocprim::half, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 10>
{};

// Based on key_type = rocprim::half, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 12>
{};

// Based on key_type = int64_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = int64_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 13>
{};

// Based on key_type = int64_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 14>
{};

// Based on key_type = int64_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int64_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = int, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<512, 31>
{};

// Based on key_type = int, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<1024, 15>
{};

// Based on key_type = int, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = int, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = short, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<512, 13>
{};

// Based on key_type = short, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<1024, 13>
{};

// Based on key_type = short, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<64, 32>
{};

// Based on key_type = short, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 32>
{};

// Based on key_type = short, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

// Based on key_type = int8_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<512, 14>
{};

// Based on key_type = int8_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<512, 24>
{};

// Based on key_type = int8_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<256, 7>
{};

// Based on key_type = int8_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 20>
{};

// Based on key_type = int8_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::unknown),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 25>
{};

// Based on key_type = double, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = double, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 13>
{};

// Based on key_type = double, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 14>
{};

// Based on key_type = double, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<128, 15>
{};

// Based on key_type = double, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = float, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = float, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<512, 31>
{};

// Based on key_type = float, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<1024, 15>
{};

// Based on key_type = float, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = float, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 31>
{};

// Based on key_type = rocprim::half, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<256, 7>
{};

// Based on key_type = rocprim::half, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<256, 9>
{};

// Based on key_type = rocprim::half, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<256, 7>
{};

// Based on key_type = rocprim::half, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 10>
{};

// Based on key_type = rocprim::half, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 12>
{};

// Based on key_type = int64_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = int64_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<256, 13>
{};

// Based on key_type = int64_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<256, 14>
{};

// Based on key_type = int64_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int64_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<256, 13>
{};

// Based on key_type = int, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<512, 31>
{};

// Based on key_type = int, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<1024, 15>
{};

// Based on key_type = int, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = int, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 15>
{};

// Based on key_type = short, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 8)
                      && (sizeof(value_type) > 4))>> : kernel_config<512, 13>
{};

// Based on key_type = short, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 4)
                      && (sizeof(value_type) > 2))>> : kernel_config<1024, 13>
{};

// Based on key_type = short, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 2)
                      && (sizeof(value_type) > 1))>> : kernel_config<64, 32>
{};

// Based on key_type = short, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1) && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 32>
{};

// Based on key_type = short, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

// Based on key_type = int8_t, value_type = int64_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 8) && (sizeof(value_type) > 4))>>
    : kernel_config<512, 14>
{};

// Based on key_type = int8_t, value_type = int
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 4) && (sizeof(value_type) > 2))>>
    : kernel_config<512, 24>
{};

// Based on key_type = int8_t, value_type = short
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 2) && (sizeof(value_type) > 1))>>
    : kernel_config<256, 7>
{};

// Based on key_type = int8_t, value_type = int8_t
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (sizeof(value_type) <= 1)
                      && (!std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 20>
{};

// Based on key_type = int8_t, value_type = empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx90a),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<1024, 25>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906), double>
    : kernel_config<256, 15>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906), float>
    : kernel_config<512, 25>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            int,
                                            double> : kernel_config<256, 13>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            int,
                                            double2> : kernel_config<256, 5>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            int,
                                            float> : kernel_config<64, 17>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906), int>
    : kernel_config<512, 25>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            int64_t,
                                            double> : kernel_config<256, 13>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            int64_t,
                                            float> : kernel_config<256, 13>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906), int64_t>
    : kernel_config<256, 15>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            int8_t,
                                            int8_t> : kernel_config<256, 20>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906), int8_t>
    : kernel_config<256, 32>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            rocprim::half> : kernel_config<256, 12>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            rocprim::half,
                                            rocprim::half> : kernel_config<512, 6>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906), short>
    : kernel_config<256, 32>
{};

template<>
struct default_radix_sort_block_sort_config<static_cast<unsigned int>(target_arch::gfx906),
                                            uint8_t,
                                            uint8_t> : kernel_config<256, 20>
{};

// Based on key_type = double, value_type = rocprim::empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx906),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = float, value_type = rocprim::empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx906),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 25>
{};

// Based on key_type = rocprim::half, value_type = rocprim::empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx906),
    key_type,
    value_type,
    std::enable_if_t<(bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 12>
{};

// Based on key_type = int64_t, value_type = rocprim::empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx906),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 8)
                      && (sizeof(key_type) > 4)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 15>
{};

// Based on key_type = int, value_type = rocprim::empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx906),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 4)
                      && (sizeof(key_type) > 2)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<512, 25>
{};

// Based on key_type = short, value_type = rocprim::empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx906),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 2)
                      && (sizeof(key_type) > 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

// Based on key_type = int8_t, value_type = rocprim::empty_type
template<class key_type, class value_type>
struct default_radix_sort_block_sort_config<
    static_cast<unsigned int>(target_arch::gfx906),
    key_type,
    value_type,
    std::enable_if_t<(!bool(rocprim::is_floating_point<key_type>::value) && (sizeof(key_type) <= 1)
                      && (std::is_same<value_type, rocprim::empty_type>::value))>>
    : kernel_config<256, 32>
{};

} // end namespace detail

END_ROCPRIM_NAMESPACE

/// @}
// end of group primitivesmodule_deviceconfigs

#endif // ROCPRIM_DEVICE_DETAIL_CONFIG_DEVICE_RADIX_SORT_BLOCK_SORT_HPP_