Implement sort
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -60,3 +60,4 @@ add_hipthrust_test("thrust.hip.copy" test_copy.cpp)
add_hipthrust_test("thrust.hip.for_each" test_for_each.cpp)
add_hipthrust_test("thrust.hip.memory" test_memory.cpp)
add_hipthrust_test("thrust.hip.transform" test_transform.cpp)
+add_hipthrust_test("thrust.hip.sort" test_sort.cpp)

ADD
test/test_sort.cpp
+// MIT License
+//
+// Copyright (c) 2019 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.
+
+#include <iostream>
+#include <type_traits>
+#include <cstdlib>
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/sort.h>
+// STREAMHPC TODO replace <thrust/detail/seq.h> with <thrust/execution_policy.h>
+#include <thrust/detail/seq.h>
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+
+
+// STREAMHPC TODO move to utils.h
+#include <algorithm>
+#include <random>
+
+template<class T>
+inline auto get_random_data(size_t size, T min, T max)
+ -> typename std::enable_if<rocprim::is_integral<T>::value, thrust::host_vector<T>>::type
+{
+ std::random_device rd;
+ std::default_random_engine gen(rd());
+ std::uniform_int_distribution<T> distribution(min, max);
+ thrust::host_vector<T> data(size);
+ std::generate(data.begin(), data.end(), [&]() { return distribution(gen); });
+ return data;
+}
+
+template<class T>
+inline auto get_random_data(size_t size, T min, T max)
+ -> typename std::enable_if<rocprim::is_floating_point<T>::value, thrust::host_vector<T>>::type
+{
+ std::random_device rd;
+ std::default_random_engine gen(rd());
+ std::uniform_real_distribution<T> distribution(min, max);
+ thrust::host_vector<T> data(size);
+ std::generate(data.begin(), data.end(), [&]() { return distribution(gen); });
+ return data;
+}
+
+template<class T>
+struct custom_compare
+{
+ __host__ __device__
+ bool operator()(const T &lhs, const T &rhs) const
+ {
+ return lhs < rhs;
+ }
+}; // end less
+
+template<
+ class Key,
+ class Item,
+ class CompareFunction = thrust::less<Key>
+>
+struct Params
+{
+ using key_type = Key;
+ using value_type = Item;
+ using compare_function = CompareFunction;
+};
+
+template<class Params>
+class SortTests : public ::testing::Test
+{
+public:
+ using key_type = typename Params::key_type;
+ using value_type = typename Params::value_type;
+ using compare_function = typename Params::compare_function;
+};
+
+typedef ::testing::Types<
+ Params<unsigned short, int, thrust::less<unsigned short> >,
+ Params<unsigned short, int, thrust::greater<unsigned short> >,
+ Params<unsigned short, int, custom_compare<unsigned short> >,
+ Params<unsigned short, double>,
+ Params<int, long long>
+> SortTestsParams;
+
+TYPED_TEST_CASE(SortTests, SortTestsParams);
+
+std::vector<size_t> get_sizes()
+{
+ std::vector<size_t> sizes = {
+ 0, 1, 2, 12, 63, 64, 211, 256, 344,
+ 1024, 2048, 5096, 34567, (1 << 17) - 1220, 1000000, (1 << 20) - 123
+ };
+ return sizes;
+}
+
+TYPED_TEST(SortTests, Sort)
+{
+ using key_type = typename TestFixture::key_type;
+ using compare_function = typename TestFixture::compare_function;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<key_type> h_keys;
+ if (std::is_floating_point<key_type>::value)
+ {
+ h_keys = get_random_data<key_type>(size, (key_type)-1000, (key_type)+1000);
+ }
+ else
+ {
+ h_keys = get_random_data<key_type>(
+ size,
+ std::numeric_limits<key_type>::min(),
+ std::numeric_limits<key_type>::max()
+ );
+ }
+
+ // Calculate expected results on host
+ thrust::host_vector<key_type> expected(h_keys);
+ thrust::sort(expected.begin(), expected.end(), compare_function());
+
+ thrust::device_vector<key_type> d_keys(h_keys);
+ thrust::sort(d_keys.begin(), d_keys.end(), compare_function());
+
+ h_keys = d_keys;
+ for (size_t i = 0; i < size; i++)
+ {
+ ASSERT_EQ(h_keys[i], expected[i]) << "where index = " << i;
+ }
+ }
+}
+
+TYPED_TEST(SortTests, SortByKey)
+{
+ using key_type = typename TestFixture::key_type;
+ using value_type = typename TestFixture::value_type;
+ using compare_function = typename TestFixture::compare_function;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ // Check if non-stable sort can be used (no equal keys with different values)
+ if (size > static_cast<size_t>(std::numeric_limits<key_type>::max())) continue;
+
+ thrust::host_vector<key_type> h_keys(size);
+ std::iota(h_keys.begin(), h_keys.end(), 0);
+ std::shuffle(
+ h_keys.begin(),
+ h_keys.end(),
+ std::default_random_engine(std::random_device{}())
+ );
+
+ thrust::host_vector<value_type> h_values(size);
+ std::iota(h_values.begin(), h_values.end(), 0);
+
+ // Calculate expected results on host
+ thrust::host_vector<key_type> expected_keys(h_keys);
+ thrust::host_vector<value_type> expected_values(h_values);
+ thrust::sort_by_key(
+ expected_keys.begin(), expected_keys.end(), expected_values.begin(),
+ compare_function());
+
+ thrust::device_vector<key_type> d_keys(h_keys);
+ thrust::device_vector<value_type> d_values(h_values);
+ thrust::sort_by_key(
+ d_keys.begin(), d_keys.end(), d_values.begin(),
+ compare_function());
+
+ h_keys = d_keys;
+ h_values = d_values;
+ for (size_t i = 0; i < size; i++)
+ {
+ ASSERT_EQ(h_keys[i], expected_keys[i]) << "where index = " << i;
+ ASSERT_EQ(h_values[i], expected_values[i]) << "where index = " << i;
+ }
+ }
+}
+
+TYPED_TEST(SortTests, StableSort)
+{
+ using key_type = typename TestFixture::key_type;
+ using compare_function = typename TestFixture::compare_function;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<key_type> h_keys;
+ if (std::is_floating_point<key_type>::value)
+ {
+ h_keys = get_random_data<key_type>(size, (key_type)-1000, (key_type)+1000);
+ }
+ else
+ {
+ h_keys = get_random_data<key_type>(
+ size,
+ std::numeric_limits<key_type>::min(),
+ std::numeric_limits<key_type>::max()
+ );
+ }
+
+ // Calculate expected results on host
+ thrust::host_vector<key_type> expected(h_keys);
+ thrust::stable_sort(expected.begin(), expected.end(), compare_function());
+
+ thrust::device_vector<key_type> d_keys(h_keys);
+ thrust::stable_sort(d_keys.begin(), d_keys.end(), compare_function());
+
+ h_keys = d_keys;
+ for (size_t i = 0; i < size; i++)
+ {
+ ASSERT_EQ(h_keys[i], expected[i]) << "where index = " << i;
+ }
+ }
+}
+
+TYPED_TEST(SortTests, StableSortByKey)
+{
+ using key_type = typename TestFixture::key_type;
+ using value_type = typename TestFixture::value_type;
+ using compare_function = typename TestFixture::compare_function;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<key_type> h_keys;
+ if (std::is_floating_point<key_type>::value)
+ {
+ h_keys = get_random_data<key_type>(size, (key_type)-1000, (key_type)+1000);
+ }
+ else
+ {
+ h_keys = get_random_data<key_type>(
+ size,
+ std::numeric_limits<key_type>::min(),
+ std::numeric_limits<key_type>::max()
+ );
+ }
+
+ thrust::host_vector<value_type> h_values(size);
+ std::iota(h_values.begin(), h_values.end(), 0);
+
+ // Calculate expected results on host
+ thrust::host_vector<key_type> expected_keys(h_keys);
+ thrust::host_vector<value_type> expected_values(h_values);
+ thrust::stable_sort_by_key(
+ expected_keys.begin(), expected_keys.end(), expected_values.begin(),
+ compare_function());
+
+ thrust::device_vector<key_type> d_keys(h_keys);
+ thrust::device_vector<value_type> d_values(h_values);
+ thrust::stable_sort_by_key(
+ d_keys.begin(), d_keys.end(), d_values.begin(),
+ compare_function());
+
+ h_keys = d_keys;
+ h_values = d_values;
+ for (size_t i = 0; i < size; i++)
+ {
+ ASSERT_EQ(h_keys[i], expected_keys[i]) << "where index = " << i;
+ ASSERT_EQ(h_values[i], expected_values[i]) << "where index = " << i;
+ }
+ }
+}

MODIFY
testing/CMakeLists.txt
testing/CMakeLists.txt
@@ -150,15 +150,15 @@ endfunction()
# add_thrust_test("set_union")
# add_thrust_test("set_union_descending")
# add_thrust_test("set_union_key_value")
-# add_thrust_test("sort_by_key")
-# add_thrust_test("sort_by_key_variable_bits")
-# add_thrust_test("sort")
+add_thrust_test("sort_by_key")
+add_thrust_test("sort_by_key_variable_bits")
+add_thrust_test("sort")
# add_thrust_test("sort_permutation_iterator")
-# add_thrust_test("sort_variable_bits")
-# add_thrust_test("stable_sort_by_key")
-# add_thrust_test("stable_sort_by_key_large")
-# add_thrust_test("stable_sort")
-# add_thrust_test("stable_sort_large")
+add_thrust_test("sort_variable_bits")
+add_thrust_test("stable_sort_by_key")
+add_thrust_test("stable_sort_by_key_large")
+add_thrust_test("stable_sort")
+add_thrust_test("stable_sort_large")
# add_thrust_test("swap_ranges")
# add_thrust_test("tabulate")
add_thrust_test("transform")

MODIFY
testing/stable_sort.cu
testing/stable_sort.cu
@@ -132,7 +132,7 @@ struct comp_mod3
comp_mod3(T * table) : table(table) {}
__host__ __device__
- bool operator()(T a, T b)
+ bool operator()(T a, T b) const
{
return table[(int) a] < table[(int) b];
}
@@ -172,4 +172,3 @@ void TestStableSortWithIndirection(void)
ASSERT_EQUAL(data[6], T(2));
}
DECLARE_VECTOR_UNITTEST(TestStableSortWithIndirection);
-

MODIFY
testing/stable_sort_large.cu
testing/stable_sort_large.cu
@@ -28,6 +28,9 @@ void TestStableSortWithLargeKeys(void)
_TestStableSortWithLargeKeys<int, 8>();
_TestStableSortWithLargeKeys<int, 16>();
_TestStableSortWithLargeKeys<int, 32>();
+// STREAM HPC investigate and fix `error: local memory limit exceeded`
+// (make block size smaller for large keys and values in rocPRIM)
+#if THRUST_DEVICE_SYSTEM != THRUST_DEVICE_SYSTEM_HIP
_TestStableSortWithLargeKeys<int, 64>();
_TestStableSortWithLargeKeys<int, 128>();
_TestStableSortWithLargeKeys<int, 256>();
@@ -38,6 +41,6 @@ void TestStableSortWithLargeKeys(void)
// _TestStableSortWithLargeKeys<int, 2048>();
// _TestStableSortWithLargeKeys<int, 4096>();
// _TestStableSortWithLargeKeys<int, 8192>();
+#endif
}
DECLARE_UNITTEST(TestStableSortWithLargeKeys);
-

ADD
thrust/system/hip/detail/sort.h
+/******************************************************************************
+ * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+#pragma once
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+#include <thrust/system/hip/detail/execution_policy.h>
+#include <thrust/system/hip/detail/par_to_seq.h>
+#include <thrust/system/hip/detail/memory_buffer.h>
+#include <thrust/sort.h>
+#include <thrust/distance.h>
+
+// rocPRIM includes
+#include <rocprim/rocprim.hpp>
+
+BEGIN_NS_THRUST
+namespace hip_rocprim {
+
+namespace __merge_sort {
+
+ template <class SORT_ITEMS>
+ struct dispatch;
+
+ // sort keys
+ template <>
+ struct dispatch<detail::false_type>
+ {
+ template <class KeysIt, class ItemsIt, class Size, class CompareOp>
+ THRUST_HIP_RUNTIME_FUNCTION static hipError_t
+ doit(void* d_temp_storage,
+ size_t& temp_storage_bytes,
+ KeysIt keys,
+ ItemsIt /*items*/,
+ Size count,
+ CompareOp compare_op,
+ hipStream_t stream,
+ bool debug_sync)
+ {
+ return rocprim::merge_sort(d_temp_storage,
+ temp_storage_bytes,
+ keys, keys,
+ static_cast<unsigned int>(count),
+ compare_op,
+ stream,
+ debug_sync);
+ }
+ };
+
+ // sort pairs
+ template <>
+ struct dispatch<detail::true_type>
+ {
+ template <class KeysIt, class ItemsIt, class Size, class CompareOp>
+ THRUST_HIP_RUNTIME_FUNCTION static hipError_t
+ doit(void* d_temp_storage,
+ size_t& temp_storage_bytes,
+ KeysIt keys,
+ ItemsIt items,
+ Size count,
+ CompareOp compare_op,
+ hipStream_t stream,
+ bool debug_sync)
+ {
+ return rocprim::merge_sort(d_temp_storage,
+ temp_storage_bytes,
+ keys, keys,
+ items, items,
+ static_cast<unsigned int>(count),
+ compare_op,
+ stream,
+ debug_sync);
+ }
+ };
+
+ template <class SORT_ITEMS,
+ class Policy,
+ class KeysIt,
+ class ItemsIt,
+ class CompareOp>
+ THRUST_HIP_RUNTIME_FUNCTION void
+ merge_sort(Policy& policy,
+ KeysIt keys_first,
+ KeysIt keys_last,
+ ItemsIt items_first,
+ CompareOp compare_op)
+ {
+ typedef typename iterator_traits<KeysIt>::difference_type size_type;
+
+ const size_type count = static_cast<size_type>(thrust::distance(keys_first, keys_last));
+
+ void* d_temp_storage = NULL;
+ size_t temp_storage_bytes = 0;
+ hipStream_t stream = hip_rocprim::stream(policy);
+ bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
+
+ hipError_t status;
+
+ status = dispatch<SORT_ITEMS>::doit(d_temp_storage,
+ temp_storage_bytes,
+ keys_first,
+ items_first,
+ count,
+ compare_op,
+ stream,
+ debug_sync);
+ hip_rocprim::throw_on_error(status, "merge_sort: failed on 1st step");
+
+ d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "merge_sort: failed to get memory buffer");
+
+ status = dispatch<SORT_ITEMS>::doit(d_temp_storage,
+ temp_storage_bytes,
+ keys_first,
+ items_first,
+ count,
+ compare_op,
+ stream,
+ debug_sync);
+ hip_rocprim::throw_on_error(status, "merge_sort: failed on 2nd step");
+
+ hip_rocprim::return_memory_buffer(policy, d_temp_storage);
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "merge_sort: failed to return memory buffer");
+ }
+} // namespace __merge_sort
+
+namespace __radix_sort {
+
+ template <class SORT_ITEMS, class Comparator>
+ struct dispatch;
+
+ // sort keys in ascending order
+ template <class K>
+ struct dispatch<detail::false_type, thrust::less<K> >
+ {
+ template <class KeysIt, class ItemsIt, class Size>
+ THRUST_HIP_RUNTIME_FUNCTION static hipError_t
+ doit(void* d_temp_storage,
+ size_t& temp_storage_bytes,
+ KeysIt keys,
+ ItemsIt /*items*/,
+ Size count,
+ hipStream_t stream,
+ bool debug_sync)
+ {
+ return rocprim::radix_sort_keys(d_temp_storage,
+ temp_storage_bytes,
+ keys, keys,
+ static_cast<unsigned int>(count),
+ 0,
+ sizeof(K) * 8,
+ stream,
+ debug_sync);
+ }
+ }; // struct dispatch -- sort keys in ascending order;
+
+ // sort keys in descending order
+ template <class K>
+ struct dispatch<detail::false_type, thrust::greater<K> >
+ {
+ template <class KeysIt, class ItemsIt, class Size>
+ THRUST_HIP_RUNTIME_FUNCTION static hipError_t
+ doit(void* d_temp_storage,
+ size_t& temp_storage_bytes,
+ KeysIt keys,
+ ItemsIt /*items*/,
+ Size count,
+ hipStream_t stream,
+ bool debug_sync)
+ {
+ return rocprim::radix_sort_keys_desc(d_temp_storage,
+ temp_storage_bytes,
+ keys, keys,
+ static_cast<unsigned int>(count),
+ 0,
+ sizeof(K) * 8,
+ stream,
+ debug_sync);
+ }
+ }; // struct dispatch -- sort keys in descending order;
+
+ // sort pairs in ascending order
+ template <class K>
+ struct dispatch<detail::true_type, thrust::less<K> >
+ {
+ template <class KeysIt, class ItemsIt, class Size>
+ THRUST_HIP_RUNTIME_FUNCTION static hipError_t
+ doit(void* d_temp_storage,
+ size_t& temp_storage_bytes,
+ KeysIt keys,
+ ItemsIt items,
+ Size count,
+ hipStream_t stream,
+ bool debug_sync)
+ {
+ return rocprim::radix_sort_pairs(d_temp_storage,
+ temp_storage_bytes,
+ keys, keys,
+ items, items,
+ static_cast<unsigned int>(count),
+ 0,
+ sizeof(K) * 8,
+ stream,
+ debug_sync);
+ }
+ }; // struct dispatch -- sort pairs in ascending order;
+
+ // sort pairs in descending order
+ template <class K>
+ struct dispatch<detail::true_type, thrust::greater<K> >
+ {
+ template <class KeysIt, class ItemsIt, class Size>
+ THRUST_HIP_RUNTIME_FUNCTION static hipError_t
+ doit(void* d_temp_storage,
+ size_t& temp_storage_bytes,
+ KeysIt keys,
+ ItemsIt items,
+ Size count,
+ hipStream_t stream,
+ bool debug_sync)
+ {
+ return rocprim::radix_sort_pairs_desc(d_temp_storage,
+ temp_storage_bytes,
+ keys, keys,
+ items, items,
+ static_cast<unsigned int>(count),
+ 0,
+ sizeof(K) * 8,
+ stream,
+ debug_sync);
+ }
+ }; // struct dispatch -- sort pairs in descending order;
+
+
+ template <class SORT_ITEMS,
+ class Policy,
+ class KeysIt,
+ class ItemsIt,
+ class CompareOp>
+ THRUST_HIP_RUNTIME_FUNCTION void
+ radix_sort(Policy& policy,
+ KeysIt keys_first,
+ KeysIt keys_last,
+ ItemsIt items_first,
+ CompareOp)
+ {
+ typedef typename iterator_traits<KeysIt>::difference_type size_type;
+
+ const size_type count = static_cast<size_type>(thrust::distance(keys_first, keys_last));
+
+ void* d_temp_storage = NULL;
+ size_t temp_storage_bytes = 0;
+ hipStream_t stream = hip_rocprim::stream(policy);
+ bool debug_sync = THRUST_HIP_DEBUG_SYNC_FLAG;
+
+ hipError_t status;
+
+ status = dispatch<SORT_ITEMS, CompareOp>::doit(d_temp_storage,
+ temp_storage_bytes,
+ keys_first,
+ items_first,
+ count,
+ stream,
+ debug_sync);
+ hip_rocprim::throw_on_error(status, "radix_sort: failed on 1st step");
+
+ d_temp_storage = hip_rocprim::get_memory_buffer(policy, temp_storage_bytes);
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "radix_sort: failed to get memory buffer");
+
+ status = dispatch<SORT_ITEMS, CompareOp>::doit(d_temp_storage,
+ temp_storage_bytes,
+ keys_first,
+ items_first,
+ count,
+ stream,
+ debug_sync);
+ hip_rocprim::throw_on_error(status, "radix_sort: failed on 2nd step");
+
+ hip_rocprim::return_memory_buffer(policy, d_temp_storage);
+ hip_rocprim::throw_on_error(hipGetLastError(),
+ "radix_sort: failed to return memory buffer");
+ }
+} // __radix_sort
+
+//---------------------------------------------------------------------
+// Smart sort picks at runtime whether to dispatch radix or merge sort
+//---------------------------------------------------------------------
+
+namespace __smart_sort {
+
+ template <class Key, class CompareOp>
+ struct can_use_primitive_sort
+ : thrust::detail::and_<
+ thrust::detail::is_arithmetic<Key>,
+ thrust::detail::or_<
+ thrust::detail::is_same<CompareOp, thrust::less<Key> >,
+ thrust::detail::is_same<CompareOp, thrust::greater<Key> > > > {};
+
+ template <class Iterator, class CompareOp>
+ struct enable_if_primitive_sort
+ : thrust::detail::enable_if<
+ can_use_primitive_sort<typename iterator_value<Iterator>::type,
+ CompareOp>::value> {};
+
+ template <class Iterator, class CompareOp>
+ struct enable_if_comparison_sort
+ : thrust::detail::disable_if<
+ can_use_primitive_sort<typename iterator_value<Iterator>::type,
+ CompareOp>::value> {};
+
+
+ template <class SORT_ITEMS,
+ class Policy,
+ class KeysIt,
+ class ItemsIt,
+ class CompareOp>
+ THRUST_HIP_RUNTIME_FUNCTION typename enable_if_comparison_sort<KeysIt, CompareOp>::type
+ smart_sort(Policy& policy,
+ KeysIt keys_first,
+ KeysIt keys_last,
+ ItemsIt items_first,
+ CompareOp compare_op)
+ {
+ __merge_sort::merge_sort<SORT_ITEMS>(policy,
+ keys_first,
+ keys_last,
+ items_first,
+ compare_op);
+ }
+
+ template <class SORT_ITEMS,
+ class Policy,
+ class KeysIt,
+ class ItemsIt,
+ class CompareOp>
+ THRUST_HIP_RUNTIME_FUNCTION typename enable_if_primitive_sort<KeysIt, CompareOp>::type
+ smart_sort(execution_policy<Policy>& policy,
+ KeysIt keys_first,
+ KeysIt keys_last,
+ ItemsIt items_first,
+ CompareOp compare_op)
+ {
+ __radix_sort::radix_sort<SORT_ITEMS>(policy,
+ keys_first,
+ keys_last,
+ items_first,
+ compare_op);
+ }
+}; // namespace __smart_sort
+
+
+//-------------------------
+// Thrust API entry points
+//-------------------------
+
+
+__thrust_exec_check_disable__
+template <class Derived, class ItemsIt, class CompareOp>
+void THRUST_HIP_FUNCTION
+stable_sort(execution_policy<Derived>& policy,
+ ItemsIt first,
+ ItemsIt last,
+ CompareOp compare_op)
+{
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND((
+ __smart_sort::smart_sort<detail::false_type, Derived, ItemsIt, ItemsIt, CompareOp>
+ ));
+#if __THRUST_HAS_HIPRT__
+ typedef typename thrust::iterator_value<ItemsIt>::type item_type;
+ __smart_sort::smart_sort<detail::false_type>(
+ policy, first, last, (item_type*)NULL, compare_op);
+#else
+ thrust::stable_sort(cvt_to_seq(derived_cast(policy)), first, last, compare_op);
+#endif
+}
+
+__thrust_exec_check_disable__
+template <class Derived, class ItemsIt, class CompareOp>
+void THRUST_HIP_FUNCTION
+sort(execution_policy<Derived>& policy,
+ ItemsIt first,
+ ItemsIt last,
+ CompareOp compare_op)
+{
+ hip_rocprim::stable_sort(policy, first, last, compare_op);
+}
+
+__thrust_exec_check_disable__
+template <class Derived,
+ class KeysIt,
+ class ValuesIt,
+ class CompareOp>
+void THRUST_HIP_FUNCTION
+stable_sort_by_key(execution_policy<Derived> &policy,
+ KeysIt keys_first,
+ KeysIt keys_last,
+ ValuesIt values,
+ CompareOp compare_op)
+{
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND((
+ __smart_sort::smart_sort<detail::true_type, Derived, KeysIt, ValuesIt, CompareOp>
+ ));
+#if __THRUST_HAS_HIPRT__
+ __smart_sort::smart_sort<detail::true_type>(
+ policy, keys_first, keys_last, values, compare_op);
+#else
+ thrust::stable_sort_by_key(
+ cvt_to_seq(derived_cast(policy)), keys_first, keys_last, values, compare_op);
+#endif
+}
+
+__thrust_exec_check_disable__
+template <class Derived, class KeysIt, class ValuesIt, class CompareOp>
+void THRUST_HIP_FUNCTION
+sort_by_key(execution_policy<Derived>& policy,
+ KeysIt keys_first,
+ KeysIt keys_last,
+ ValuesIt values,
+ CompareOp compare_op)
+{
+ hip_rocprim::stable_sort_by_key(policy, keys_first, keys_last, values, compare_op);
+}
+
+// API with default comparator
+
+template <class Derived, class ItemsIt>
+void THRUST_HIP_FUNCTION
+sort(execution_policy<Derived>& policy,
+ ItemsIt first,
+ ItemsIt last)
+{
+ typedef typename thrust::iterator_value<ItemsIt>::type item_type;
+ hip_rocprim::sort(policy, first, last, less<item_type>());
+}
+
+template <class Derived, class ItemsIt>
+void THRUST_HIP_FUNCTION
+stable_sort(execution_policy<Derived>& policy,
+ ItemsIt first,
+ ItemsIt last)
+{
+ typedef typename thrust::iterator_value<ItemsIt>::type item_type;
+ hip_rocprim::stable_sort(policy, first, last, less<item_type>());
+}
+
+template <class Derived, class KeysIt, class ValuesIt>
+void THRUST_HIP_FUNCTION
+sort_by_key(execution_policy<Derived>& policy,
+ KeysIt keys_first,
+ KeysIt keys_last,
+ ValuesIt values)
+{
+ typedef typename thrust::iterator_value<KeysIt>::type key_type;
+ hip_rocprim::sort_by_key(policy, keys_first, keys_last, values, less<key_type>());
+}
+
+template <class Derived, class KeysIt, class ValuesIt>
+void THRUST_HIP_FUNCTION
+stable_sort_by_key(execution_policy<Derived>& policy,
+ KeysIt keys_first,
+ KeysIt keys_last,
+ ValuesIt values)
+{
+ typedef typename thrust::iterator_value<KeysIt>::type key_type;
+ hip_rocprim::stable_sort_by_key(policy, keys_first, keys_last, values, less<key_type>());
+}
+
+
+} // namespace hip_rocprim
+END_NS_THRUST
+#endif

MODIFY
thrust/system/hip/execution_policy.h
thrust/system/hip/execution_policy.h
// #include <thrust/system/hip/detail/merge.h>
// #include <thrust/system/hip/detail/scan_by_key.h>
// #include <thrust/system/hip/detail/set_operations.h>
-// #include <thrust/system/hip/detail/sort.h>
+#include <thrust/system/hip/detail/sort.h>
// // work in progress
-

Implement sort