Implemented stable sort tests
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -86,6 +86,10 @@ add_hipthrust_test("thrust.hip.remove" test_remove.cpp)
add_hipthrust_test("thrust.hip.replace" test_replace.cpp)
add_hipthrust_test("thrust.hip.reverse_iterator" test_reverse_iterator.cpp)
add_hipthrust_test("thrust.hip.scatter" test_scatter.cpp)
+add_hipthrust_test("thrust.hip.stable_sort" test_stable_sort.cpp)
+add_hipthrust_test("thrust.hip.stable_sort_by_key" test_stable_sort_by_key.cpp)
+add_hipthrust_test("thrust.hip.stable_sort_by_key_large" test_stable_sort_by_key_large.cpp)
+add_hipthrust_test("thrust.hip.stable_sort_large" test_stable_sort_large.cpp)
add_hipthrust_test("thrust.hip.sort" test_sort.cpp)
add_hipthrust_test("thrust.hip.tabulate" test_tabulate.cpp)
add_hipthrust_test("thrust.hip.transform" test_transform.cpp)

ADD
test/test_stable_sort.cpp
+// MIT License
+//
+// Copyright (c) 2018 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.
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/sort.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/retag.h>
+
+#include "test_utils.hpp"
+
+template<
+ class Input
+>
+struct Params
+{
+ using input_type = Input;
+};
+
+template<class Params>
+class StableSortTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<short>,
+ Params<int>,
+ Params<long long>,
+ Params<thrust::detail::int8_t>,
+ Params<thrust::detail::int16_t>,
+ Params<thrust::detail::int32_t>,
+ Params<thrust::detail::int64_t>
+> StableSortTestsParams;
+
+TYPED_TEST_CASE(StableSortTests, StableSortTestsParams);
+
+template<class Params>
+class StableSortVectorTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<thrust::host_vector<short>>,
+ Params<thrust::host_vector<int>>,
+ Params<thrust::host_vector<long long>>,
+ Params<thrust::host_vector<unsigned short>>,
+ Params<thrust::host_vector<unsigned int>>,
+ Params<thrust::host_vector<unsigned long long>>,
+ Params<thrust::device_vector<short>>,
+ Params<thrust::device_vector<int>>,
+ Params<thrust::device_vector<long long>>,
+ Params<thrust::device_vector<unsigned short>>,
+ Params<thrust::device_vector<unsigned int>>,
+ Params<thrust::device_vector<unsigned long long>>
+> StableSortVectorTestsParams;
+
+TYPED_TEST_CASE(StableSortVectorTests, StableSortVectorTestsParams);
+
+template<typename RandomAccessIterator>
+void stable_sort(my_system &system, RandomAccessIterator, RandomAccessIterator)
+{
+ system.validate_dispatch();
+}
+
+TEST(StableSortTests, TestStableSortDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::stable_sort(sys, vec.begin(), vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename RandomAccessIterator>
+void stable_sort(my_tag, RandomAccessIterator first, RandomAccessIterator)
+{
+ *first = 13;
+}
+
+TEST(StableSortTests, TestStableSortDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::stable_sort(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()));
+
+ ASSERT_EQ(13, vec.front());
+}
+
+template <typename T>
+struct less_div_10
+{
+ __host__ __device__ bool operator()(const T &lhs, const T &rhs) const {return ((int) lhs) / 10 < ((int) rhs) / 10;}
+};
+
+template <class Vector>
+void InitializeSimpleStableKeySortTest(Vector& unsorted_keys, Vector& sorted_keys)
+{
+ unsorted_keys.resize(9);
+ unsorted_keys[0] = 25;
+ unsorted_keys[1] = 14;
+ unsorted_keys[2] = 35;
+ unsorted_keys[3] = 16;
+ unsorted_keys[4] = 26;
+ unsorted_keys[5] = 34;
+ unsorted_keys[6] = 36;
+ unsorted_keys[7] = 24;
+ unsorted_keys[8] = 15;
+
+ sorted_keys.resize(9);
+ sorted_keys[0] = 14;
+ sorted_keys[1] = 16;
+ sorted_keys[2] = 15;
+ sorted_keys[3] = 25;
+ sorted_keys[4] = 26;
+ sorted_keys[5] = 24;
+ sorted_keys[6] = 35;
+ sorted_keys[7] = 34;
+ sorted_keys[8] = 36;
+}
+
+TYPED_TEST(StableSortVectorTests, TestStableSortSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector unsorted_keys;
+ Vector sorted_keys;
+
+ InitializeSimpleStableKeySortTest(unsorted_keys, sorted_keys);
+
+ thrust::stable_sort(unsorted_keys.begin(), unsorted_keys.end(), less_div_10<T>());
+
+ ASSERT_EQ(unsorted_keys, sorted_keys);
+}
+
+TYPED_TEST(StableSortTests, TestStableSort)
+{
+ using T = typename TestFixture::input_type;
+
+ for(auto size : get_sizes())
+ {
+ thrust::host_vector<T> h_data = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+ thrust::device_vector<T> d_data = h_data;
+
+ thrust::stable_sort(h_data.begin(), h_data.end(), less_div_10<T>());
+ thrust::stable_sort(d_data.begin(), d_data.end(), less_div_10<T>());
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+template <typename T>
+struct comp_mod3
+{
+ T * table;
+
+ comp_mod3(T * table) : table(table) {}
+
+ __host__ __device__
+ bool operator()(T a, T b) const
+ {
+ return table[(int) a] < table[(int) b];
+ }
+};
+
+TYPED_TEST(StableSortVectorTests, TestStableSortWithIndirection)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(7);
+ data[0] = T(1);
+ data[1] = T(3);
+ data[2] = T(5);
+ data[3] = T(3);
+ data[4] = T(0);
+ data[5] = T(2);
+ data[6] = T(1);
+
+ Vector table(6);
+ table[0] = T(0);
+ table[1] = T(1);
+ table[2] = T(2);
+ table[3] = T(0);
+ table[4] = T(1);
+ table[5] = T(2);
+
+ thrust::stable_sort(data.begin(), data.end(), comp_mod3<T>(thrust::raw_pointer_cast(&table[0])));
+
+ ASSERT_EQ(data[0], T(3));
+ ASSERT_EQ(data[1], T(3));
+ ASSERT_EQ(data[2], T(0));
+ ASSERT_EQ(data[3], T(1));
+ ASSERT_EQ(data[4], T(1));
+ ASSERT_EQ(data[5], T(5));
+ ASSERT_EQ(data[6], T(2));
+}

ADD
test/test_stable_sort_by_key.cpp
+// MIT License
+//
+// Copyright (c) 2018 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.
+
+// Google Test
+#include <gtest/gtest.h>
+
+// Thrust
+#include <thrust/sort.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/retag.h>
+
+#include "test_utils.hpp"
+
+template<
+ class Input,
+ class Output = Input
+>
+struct Params
+{
+ using input_type = Input;
+ using output_type = Output;
+};
+
+template<class Params>
+class StableSortByKeyTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<short>,
+ Params<int>,
+ Params<long long>,
+ Params<thrust::detail::int8_t>,
+ Params<thrust::detail::int16_t>,
+ Params<thrust::detail::int32_t>,
+ Params<thrust::detail::int64_t>
+> StableSortByKeyTestsParams;
+
+TYPED_TEST_CASE(StableSortByKeyTests, StableSortByKeyTestsParams);
+
+template<class Params>
+class StableSortByKeyVectorTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<thrust::host_vector<short>>,
+ Params<thrust::host_vector<int>>,
+ Params<thrust::host_vector<long long>>,
+ Params<thrust::host_vector<unsigned short>>,
+ Params<thrust::host_vector<unsigned int>>,
+ Params<thrust::host_vector<unsigned long long>>,
+ Params<thrust::device_vector<short>>,
+ Params<thrust::device_vector<int>>,
+ Params<thrust::device_vector<long long>>,
+ Params<thrust::device_vector<unsigned short>>,
+ Params<thrust::device_vector<unsigned int>>,
+ Params<thrust::device_vector<unsigned long long>>
+> StableSortByKeyVectorTestsParams;
+
+TYPED_TEST_CASE(StableSortByKeyVectorTests, StableSortByKeyVectorTestsParams);
+
+template<typename RandomAccessIterator1, typename RandomAccessIterator2>
+void stable_sort_by_key(my_system &system, RandomAccessIterator1, RandomAccessIterator1, RandomAccessIterator2)
+{
+ system.validate_dispatch();
+}
+
+TEST(StableSortByKeyTests, TestStableSortByKeyDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::stable_sort_by_key(sys, vec.begin(), vec.begin(), vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename RandomAccessIterator1, typename RandomAccessIterator2>
+void stable_sort_by_key(my_tag, RandomAccessIterator1 keys_first, RandomAccessIterator1, RandomAccessIterator2)
+{
+ *keys_first = 13;
+}
+
+TEST(StableSortByKeyTests, TestStableSortByKeyDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::stable_sort_by_key(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()));
+
+ ASSERT_EQ(13, vec.front());
+}
+
+template <typename T>
+struct less_div_10
+{
+ __host__ __device__ bool operator()(const T &lhs, const T &rhs) const {return ((int) lhs) / 10 < ((int) rhs) / 10;}
+};
+
+template <class Vector>
+void InitializeSimpleStableKeyValueSortTest(Vector& unsorted_keys, Vector& unsorted_values,
+ Vector& sorted_keys, Vector& sorted_values)
+{
+ unsorted_keys.resize(9);
+ unsorted_values.resize(9);
+ unsorted_keys[0] = 25; unsorted_values[0] = 0;
+ unsorted_keys[1] = 14; unsorted_values[1] = 1;
+ unsorted_keys[2] = 35; unsorted_values[2] = 2;
+ unsorted_keys[3] = 16; unsorted_values[3] = 3;
+ unsorted_keys[4] = 26; unsorted_values[4] = 4;
+ unsorted_keys[5] = 34; unsorted_values[5] = 5;
+ unsorted_keys[6] = 36; unsorted_values[6] = 6;
+ unsorted_keys[7] = 24; unsorted_values[7] = 7;
+ unsorted_keys[8] = 15; unsorted_values[8] = 8;
+
+ sorted_keys.resize(9);
+ sorted_values.resize(9);
+ sorted_keys[0] = 14; sorted_values[0] = 1;
+ sorted_keys[1] = 16; sorted_values[1] = 3;
+ sorted_keys[2] = 15; sorted_values[2] = 8;
+ sorted_keys[3] = 25; sorted_values[3] = 0;
+ sorted_keys[4] = 26; sorted_values[4] = 4;
+ sorted_keys[5] = 24; sorted_values[5] = 7;
+ sorted_keys[6] = 35; sorted_values[6] = 2;
+ sorted_keys[7] = 34; sorted_values[7] = 5;
+ sorted_keys[8] = 36; sorted_values[8] = 6;
+}
+
+TYPED_TEST(StableSortByKeyVectorTests, TestStableSortByKeySimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector unsorted_keys, unsorted_values;
+ Vector sorted_keys, sorted_values;
+
+ InitializeSimpleStableKeyValueSortTest(unsorted_keys, unsorted_values, sorted_keys, sorted_values);
+
+ thrust::stable_sort_by_key(unsorted_keys.begin(), unsorted_keys.end(), unsorted_values.begin(), less_div_10<T>());
+
+ ASSERT_EQ(unsorted_keys, sorted_keys);
+ ASSERT_EQ(unsorted_values, sorted_values);
+}
+
+TYPED_TEST(StableSortByKeyVectorTests, TestStableSortByKey)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ for(auto size : get_sizes())
+ {
+ thrust::host_vector<T> h_keys = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+ thrust::device_vector<T> d_keys = h_keys;
+
+ thrust::host_vector<T> h_values = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+ thrust::device_vector<T> d_values = h_values;
+
+ thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
+ thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
+
+ ASSERT_EQ(h_keys, d_keys);
+ ASSERT_EQ(h_values, d_values);
+ }
+}

ADD
test/test_stable_sort_by_key_large.cpp
+// MIT License
+//
+// Copyright (c) 2018 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 <thrust/sort.h>
+#include <thrust/functional.h>
+
+#include "test_utils.hpp"
+#include "test_assertions.hpp"
+
+template <typename T>
+struct less_div_10
+{
+ __host__ __device__ bool operator()(const T &lhs, const T &rhs) const {return ((int) lhs) / 10 < ((int) rhs) / 10;}
+};
+
+template <typename T>
+struct greater_div_10
+{
+ __host__ __device__ bool operator()(const T &lhs, const T &rhs) const {return ((int) lhs) / 10 > ((int) rhs) / 10;}
+};
+
+
+template <typename T, unsigned int N>
+void _TestStableSortByKeyWithLargeKeys(void)
+{
+ size_t n = (128 * 1024) / sizeof(FixedVector<T,N>);
+
+ thrust::host_vector< FixedVector<T,N> > h_keys(n);
+ thrust::host_vector< unsigned int > h_vals(n);
+
+ for(size_t i = 0; i < n; i++)
+ {
+ h_keys[i] = FixedVector<T,N>(rand());
+ h_vals[i] = i;
+ }
+
+ thrust::device_vector< FixedVector<T,N> > d_keys = h_keys;
+ thrust::device_vector< unsigned int > d_vals = h_vals;
+
+ thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_vals.begin());
+ thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_vals.begin());
+
+ ASSERT_EQ_QUIET(h_keys, d_keys);
+ ASSERT_EQ_QUIET(h_vals, d_vals);
+}
+
+
+TEST(StableSortByKeyLargeTests, TestStableSortByKeyWithLargeKeys)
+{
+ _TestStableSortByKeyWithLargeKeys<int, 4>();
+ _TestStableSortByKeyWithLargeKeys<int, 8>();
+ _TestStableSortByKeyWithLargeKeys<int, 16>();
+
+// XXX these take too long to compile
+// _TestStableSortByKeyWithLargeKeys<int, 32>();
+// _TestStableSortByKeyWithLargeKeys<int, 64>();
+// _TestStableSortByKeyWithLargeKeys<int, 128>();
+// _TestStableSortByKeyWithLargeKeys<int, 256>();
+// _TestStableSortByKeyWithLargeKeys<int, 512>();
+// _TestStableSortByKeyWithLargeKeys<int, 1024>();
+// _TestStableSortByKeyWithLargeKeys<int, 2048>();
+// _TestStableSortByKeyWithLargeKeys<int, 4096>();
+// _TestStableSortByKeyWithLargeKeys<int, 8192>();
+}
+
+template <typename T, unsigned int N>
+void _TestStableSortByKeyWithLargeValues(void)
+{
+ size_t n = (128 * 1024) / sizeof(FixedVector<T,N>);
+
+ thrust::host_vector< unsigned int > h_keys(n);
+ thrust::host_vector< FixedVector<T,N> > h_vals(n);
+
+ for(size_t i = 0; i < n; i++)
+ {
+ h_keys[i] = rand();
+ h_vals[i] = FixedVector<T,N>(i);
+ }
+
+ thrust::device_vector< unsigned int > d_keys = h_keys;
+ thrust::device_vector< FixedVector<T,N> > d_vals = h_vals;
+
+ thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_vals.begin());
+ thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_vals.begin());
+
+ ASSERT_EQ_QUIET(h_keys, d_keys);
+ ASSERT_EQ_QUIET(h_vals, d_vals);
+
+ // so cuda::stable_merge_sort_by_key() is called
+ thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_vals.begin(), greater_div_10<unsigned int>());
+ thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_vals.begin(), greater_div_10<unsigned int>());
+
+ ASSERT_EQ_QUIET(h_keys, d_keys);
+ ASSERT_EQ_QUIET(h_vals, d_vals);
+}
+
+TEST(StableSortByKeyLargeTests, TestStableSortByKeyWithLargeValues)
+{
+ _TestStableSortByKeyWithLargeValues<int, 4>();
+ _TestStableSortByKeyWithLargeValues<int, 8>();
+ _TestStableSortByKeyWithLargeValues<int, 16>();
+
+// XXX these take too long to compile
+// _TestStableSortByKeyWithLargeValues<int, 32>();
+// _TestStableSortByKeyWithLargeValues<int, 64>();
+// _TestStableSortByKeyWithLargeValues<int, 128>();
+// _TestStableSortByKeyWithLargeValues<int, 256>();
+// _TestStableSortByKeyWithLargeValues<int, 512>();
+// _TestStableSortByKeyWithLargeValues<int, 1024>();
+// _TestStableSortByKeyWithLargeValues<int, 2048>();
+// _TestStableSortByKeyWithLargeValues<int, 4096>();
+// _TestStableSortByKeyWithLargeValues<int, 8192>();
+}
+
+
+template <typename T, unsigned int N>
+void _TestStableSortByKeyWithLargeKeysAndValues(void)
+{
+ size_t n = (128 * 1024) / sizeof(FixedVector<T,N>);
+
+ thrust::host_vector< FixedVector<T,N> > h_keys(n);
+ thrust::host_vector< FixedVector<T,N> > h_vals(n);
+
+ for(size_t i = 0; i < n; i++)
+ {
+ h_keys[i] = FixedVector<T,N>(rand());
+ h_vals[i] = FixedVector<T,N>(i);
+ }
+
+ thrust::device_vector< FixedVector<T,N> > d_keys = h_keys;
+ thrust::device_vector< FixedVector<T,N> > d_vals = h_vals;
+
+ thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_vals.begin());
+ thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_vals.begin());
+
+ ASSERT_EQ_QUIET(h_keys, d_keys);
+ ASSERT_EQ_QUIET(h_vals, d_vals);
+}
+
+
+TEST(StableSortByKeyLargeTests, TestStableSortByKeyWithLargeKeysAndValues)
+{
+ _TestStableSortByKeyWithLargeKeysAndValues<int, 4>();
+ _TestStableSortByKeyWithLargeKeysAndValues<int, 8>();
+ _TestStableSortByKeyWithLargeKeysAndValues<int, 16>();
+
+// XXX these take too long to compile
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 32>();
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 64>();
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 128>();
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 256>();
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 512>();
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 1024>();
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 2048>();
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 4096>();
+// _TestStableSortByKeyWithLargeKeysAndValues<int, 8192>();
+}

ADD
test/test_stable_sort_large.cpp
+// MIT License
+//
+// Copyright (c) 2018 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 <thrust/sort.h>
+#include <thrust/functional.h>
+
+#include "test_utils.hpp"
+#include "test_assertions.hpp"
+
+template <typename T, unsigned int N>
+void _TestStableSortWithLargeKeys(void)
+{
+ size_t n = (128 * 1024) / sizeof(FixedVector<T,N>);
+
+ thrust::host_vector< FixedVector<T,N> > h_keys(n);
+
+ for(size_t i = 0; i < n; i++)
+ h_keys[i] = FixedVector<T,N>(rand());
+
+ thrust::device_vector< FixedVector<T,N> > d_keys = h_keys;
+
+ thrust::stable_sort(h_keys.begin(), h_keys.end());
+ thrust::stable_sort(d_keys.begin(), d_keys.end());
+
+ ASSERT_EQ_QUIET(h_keys, d_keys);
+}
+
+TEST(StableSortLargeTests, TestStableSortWithLargeKeys)
+{
+ _TestStableSortWithLargeKeys<int, 1>();
+ _TestStableSortWithLargeKeys<int, 2>();
+ _TestStableSortWithLargeKeys<int, 4>();
+ _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>();
+ _TestStableSortWithLargeKeys<int, 512>();
+ _TestStableSortWithLargeKeys<int, 1024>();
+
+// XXX these take too long to compile
+// _TestStableSortWithLargeKeys<int, 2048>();
+// _TestStableSortWithLargeKeys<int, 4096>();
+// _TestStableSortWithLargeKeys<int, 8192>();
+#endif
+}

Implemented stable sort tests