warp_sort.hpp

// Copyright (c) 2017-2021 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_WARP_WARP_SORT_HPP_
#define ROCPRIM_WARP_WARP_SORT_HPP_

#include <type_traits>

#include "../config.hpp"
#include "../detail/various.hpp"

#include "../intrinsics.hpp"
#include "../functional.hpp"

#include "detail/warp_sort_shuffle.hpp"


BEGIN_ROCPRIM_NAMESPACE

template<
    class Key,
    unsigned int WarpSize = device_warp_size(),
    class Value = empty_type
>
class warp_sort : detail::warp_sort_shuffle<Key, WarpSize, Value>
{
    typedef typename detail::warp_sort_shuffle<Key, WarpSize, Value> base_type;

    // Check if WarpSize is valid for the targets
    static_assert(WarpSize <= ROCPRIM_MAX_WARP_SIZE, "WarpSize can't be greater than hardware warp size.");

public:
    typedef typename base_type::storage_type storage_type;

    template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key& thread_key,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        base_type::sort(thread_key, compare_function);
    }

    template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key& ,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        (void) compare_function; // disables unused parameter warning
        ROCPRIM_PRINT_ERROR_ONCE("Specified warp size exceeds current hardware supported warp size. Aborting warp sort.");
        return;
    }

    template<
        unsigned int ItemsPerThread,
        class BinaryFunction = ::rocprim::less<Key>,
        unsigned int FunctionWarpSize = WarpSize
    >
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key (&thread_keys)[ItemsPerThread],
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        base_type::sort(thread_keys, compare_function);
    }

    template<
        unsigned int ItemsPerThread,
        class BinaryFunction = ::rocprim::less<Key>,
        unsigned int FunctionWarpSize = WarpSize
    >
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key (&thread_keys)[ItemsPerThread],
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        (void) thread_keys;      // disables unused parameter warning
        (void) compare_function; // disables unused parameter warning
        ROCPRIM_PRINT_ERROR_ONCE("Specified warp size exceeds current hardware supported warp size. Aborting warp sort.");
        return;
    }

    template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key& thread_key,
              storage_type& storage,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        base_type::sort(
            thread_key, storage, compare_function
        );
    }

    template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key& ,
              storage_type& ,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        (void) compare_function; // disables unused parameter warning
        ROCPRIM_PRINT_ERROR_ONCE("Specified warp size exceeds current hardware supported warp size. Aborting warp sort.");
        return;
    }


    template<
        unsigned int ItemsPerThread,
        class BinaryFunction = ::rocprim::less<Key>,
        unsigned int FunctionWarpSize = WarpSize
    >
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key (&thread_keys)[ItemsPerThread],
              storage_type& storage,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        base_type::sort(
            thread_keys, storage, compare_function
        );
    }

    template<
        unsigned int ItemsPerThread,
        class BinaryFunction = ::rocprim::less<Key>,
        unsigned int FunctionWarpSize = WarpSize
    >
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key (&thread_keys)[ItemsPerThread],
              storage_type& ,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        (void) thread_keys;      // disables unused parameter warning
        (void) compare_function; // disables unused parameter warning
        ROCPRIM_PRINT_ERROR_ONCE("Specified warp size exceeds current hardware supported warp size. Aborting warp sort.");
        return;
    }

    template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key& thread_key,
              Value& thread_value,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        base_type::sort(
            thread_key, thread_value, compare_function
        );
    }

    template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key& ,
              Value& ,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        (void) compare_function; // disables unused parameter warning
        ROCPRIM_PRINT_ERROR_ONCE("Specified warp size exceeds current hardware supported warp size. Aborting warp sort.");
        return;
    }

    template<
        unsigned int ItemsPerThread,
        class BinaryFunction = ::rocprim::less<Key>,
        unsigned int FunctionWarpSize = WarpSize
    >
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key (&thread_keys)[ItemsPerThread],
              Value (&thread_values)[ItemsPerThread],
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        base_type::sort(
            thread_keys, thread_values, compare_function
        );
    }

    template<
        unsigned int ItemsPerThread,
        class BinaryFunction = ::rocprim::less<Key>,
        unsigned int FunctionWarpSize = WarpSize
    >
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key (&thread_keys)[ItemsPerThread],
              Value (&thread_values)[ItemsPerThread],
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        (void) thread_keys;      // disables unused parameter warning
        (void) thread_values;    // disables unused parameter warning
        (void) compare_function; // disables unused parameter warning
        ROCPRIM_PRINT_ERROR_ONCE("Specified warp size exceeds current hardware supported warp size. Aborting warp sort.");
        return;
    }

    template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key& thread_key,
              Value& thread_value,
              storage_type& storage,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        base_type::sort(
            thread_key, thread_value, storage, compare_function
        );
    }

    template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key& ,
              Value& ,
              storage_type& ,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        (void) compare_function; // disables unused parameter warning
        ROCPRIM_PRINT_ERROR_ONCE("Specified warp size exceeds current hardware supported warp size. Aborting warp sort.");
        return;
    }


    template<
        unsigned int ItemsPerThread,
        class BinaryFunction = ::rocprim::less<Key>,
        unsigned int FunctionWarpSize = WarpSize
    >
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key (&thread_keys)[ItemsPerThread],
              Value (&thread_values)[ItemsPerThread],
              storage_type& storage,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        base_type::sort(
            thread_keys, thread_values, storage, compare_function
        );
    }

    template<
        unsigned int ItemsPerThread,
        class BinaryFunction = ::rocprim::less<Key>,
        unsigned int FunctionWarpSize = WarpSize
    >
    ROCPRIM_DEVICE ROCPRIM_INLINE
    auto sort(Key (&thread_keys)[ItemsPerThread],
              Value (&thread_values)[ItemsPerThread],
              storage_type& ,
              BinaryFunction compare_function = BinaryFunction())
        -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type
    {
        (void) thread_keys;      // disables unused parameter warning
        (void) thread_values;    // disables unused parameter warning
        (void) compare_function; // disables unused parameter warning
        ROCPRIM_PRINT_ERROR_ONCE("Specified warp size exceeds current hardware supported warp size. Aborting warp sort.");
        return;
    }
};

END_ROCPRIM_NAMESPACE

// end of group warpmodule

#endif // ROCPRIM_WARP_WARP_SORT_HPP_