Implemented sort tests
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -92,6 +92,9 @@ 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.sort_by_key" test_sort_by_key.cpp)
+add_hipthrust_test("thrust.hip.sort_by_key_variable_bits" test_sort_by_key_variable_bits.cpp)
+add_hipthrust_test("thrust.hip.sort_permutation_iterator" test_sort_permutation_iterator.cpp)
add_hipthrust_test("thrust.hip.tabulate" test_tabulate.cpp)
add_hipthrust_test("thrust.hip.transform" test_transform.cpp)
add_hipthrust_test("thrust.hip.transform_iterator" test_transform_iterator.cpp)

MODIFY
test/test_sort.cpp
test/test_sort.cpp
#include <thrust/detail/seq.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/retag.h>
#include "test_utils.hpp"
@@ -65,7 +67,39 @@ typedef ::testing::Types<
TYPED_TEST_CASE(SortTests, SortTestsParams);
+template< class InputType >
+struct ParamsVector
+{
+ using input_type = InputType;
+};
+
+template<class ParamsVector>
+class SortVector : public ::testing::Test
+{
+public:
+ using input_type = typename ParamsVector::input_type;
+};
+typedef ::testing::Types<
+ ParamsVector<thrust::host_vector<short>>,
+ ParamsVector<thrust::host_vector<int>>,
+ ParamsVector<thrust::host_vector<long long>>,
+ ParamsVector<thrust::host_vector<unsigned short>>,
+ ParamsVector<thrust::host_vector<unsigned int>>,
+ ParamsVector<thrust::host_vector<unsigned long long>>,
+ ParamsVector<thrust::host_vector<float>>,
+ ParamsVector<thrust::host_vector<double>>,
+ ParamsVector<thrust::device_vector<short>>,
+ ParamsVector<thrust::device_vector<int>>,
+ ParamsVector<thrust::device_vector<long long>>,
+ ParamsVector<thrust::device_vector<unsigned short>>,
+ ParamsVector<thrust::device_vector<unsigned int>>,
+ ParamsVector<thrust::device_vector<unsigned long long>>,
+ ParamsVector<thrust::device_vector<float>>,
+ ParamsVector<thrust::device_vector<double>>
+> SortVectorParams;
+
+TYPED_TEST_CASE(SortVector, SortVectorParams);
TYPED_TEST(SortTests, Sort)
{
@@ -239,3 +273,139 @@ TYPED_TEST(SortTests, StableSortByKey)
}
}
}
+
+template<typename RandomAccessIterator>
+void sort(my_system &system, RandomAccessIterator, RandomAccessIterator)
+{
+ system.validate_dispatch();
+}
+
+TEST(SortTests, TestSortDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::sort(sys, vec.begin(), vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename RandomAccessIterator>
+void sort(my_tag, RandomAccessIterator first, RandomAccessIterator)
+{
+ *first = 13;
+}
+
+TEST(SortTests, TestSortDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::sort(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()));
+
+ ASSERT_EQ(13, vec.front());
+}
+
+template <class Vector>
+void InitializeSimpleKeySortTest(Vector& unsorted_keys, Vector& sorted_keys)
+{
+ using T = typename Vector::value_type;
+
+ unsorted_keys.resize(7);
+ unsorted_keys[0] = T(1);
+ unsorted_keys[1] = T(3);
+ unsorted_keys[2] = T(6);
+ unsorted_keys[3] = T(5);
+ unsorted_keys[4] = T(2);
+ unsorted_keys[5] = T(0);
+ unsorted_keys[6] = T(4);
+
+ sorted_keys.resize(7);
+ sorted_keys[0] = T(0);
+ sorted_keys[1] = T(1);
+ sorted_keys[2] = T(2);
+ sorted_keys[3] = T(3);
+ sorted_keys[4] = T(4);
+ sorted_keys[5] = T(5);
+ sorted_keys[6] = T(6);
+}
+
+TYPED_TEST(SortVector, TestSortSimple)
+{
+ using Vector = typename TestFixture::input_type;
+
+ Vector unsorted_keys;
+ Vector sorted_keys;
+
+ InitializeSimpleKeySortTest(unsorted_keys, sorted_keys);
+
+ thrust::sort(unsorted_keys.begin(), unsorted_keys.end());
+
+ ASSERT_EQ(unsorted_keys, sorted_keys);
+}
+
+TYPED_TEST(SortVector, TestSortAscendingKey)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_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::sort(h_data.begin(), h_data.end(), thrust::less<T>());
+ thrust::sort(d_data.begin(), d_data.end(), thrust::less<T>());
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TEST(SortTests, TestSortDescendingKey)
+{
+ const size_t size = 10027;
+
+ thrust::host_vector<int> h_data = get_random_data<int>(size,
+ std::numeric_limits<int>::min(),
+ std::numeric_limits<int>::max());
+ thrust::device_vector<int> d_data = h_data;
+
+ thrust::sort(h_data.begin(), h_data.end(), thrust::greater<int>());
+ thrust::sort(d_data.begin(), d_data.end(), thrust::greater<int>());
+
+ ASSERT_EQ(h_data, d_data);
+}
+
+TEST(SortTests, TestSortBool)
+{
+ const size_t size = 10027;
+
+ thrust::host_vector<bool> h_data = get_random_data<bool>(size,
+ std::numeric_limits<bool>::min(),
+ std::numeric_limits<bool>::max());
+
+ thrust::device_vector<bool> d_data = h_data;
+
+ thrust::sort(h_data.begin(), h_data.end());
+ thrust::sort(d_data.begin(), d_data.end());
+
+ ASSERT_EQ(h_data, d_data);
+}
+
+TEST(SortTests, TestSortBoolDescending)
+{
+ const size_t size = 10027;
+
+ thrust::host_vector<bool> h_data = get_random_data<bool>(size,
+ std::numeric_limits<bool>::min(),
+ std::numeric_limits<bool>::max());
+
+ thrust::device_vector<bool> d_data = h_data;
+
+ thrust::sort(h_data.begin(), h_data.end(), thrust::greater<bool>());
+ thrust::sort(d_data.begin(), d_data.end(), thrust::greater<bool>());
+
+ ASSERT_EQ(h_data, d_data);
+}

ADD
test/test_sort_by_key.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.
+
+// Google Test
+#include <gtest/gtest.h>
+
+#include <thrust/sort.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/retag.h>
+
+#include "test_utils.hpp"
+
+template< class InputType >
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class SortByKeyTests : 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::host_vector<float>>,
+ Params<thrust::host_vector<double>>,
+ 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>>,
+ Params<thrust::device_vector<float>>,
+ Params<thrust::device_vector<double>>
+> SortByKeyTestsParams;
+
+TYPED_TEST_CASE(SortByKeyTests, SortByKeyTestsParams);
+
+template<typename RandomAccessIterator1,
+ typename RandomAccessIterator2>
+void sort_by_key(my_system &system, RandomAccessIterator1, RandomAccessIterator1, RandomAccessIterator2)
+{
+ system.validate_dispatch();
+}
+
+
+TEST(SortByKeyTests, TestSortByKeyDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::sort_by_key(sys, vec.begin(), vec.begin(), vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename RandomAccessIterator1,
+ typename RandomAccessIterator2>
+void sort_by_key(my_tag, RandomAccessIterator1 keys_first, RandomAccessIterator1, RandomAccessIterator2)
+{
+ *keys_first = 13;
+}
+
+TEST(SortByKeyTests, TestSortByKeyDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::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 <class Vector>
+void InitializeSimpleKeyValueSortTest(Vector& unsorted_keys, Vector& unsorted_values,
+ Vector& sorted_keys, Vector& sorted_values)
+{
+ using T = typename Vector::value_type;
+
+ unsorted_keys.resize(7);
+ unsorted_values.resize(7);
+ unsorted_keys[0] = T(1); unsorted_values[0] = T(0);
+ unsorted_keys[1] = T(3); unsorted_values[1] = T(1);
+ unsorted_keys[2] = T(6); unsorted_values[2] = T(2);
+ unsorted_keys[3] = T(5); unsorted_values[3] = T(3);
+ unsorted_keys[4] = T(2); unsorted_values[4] = T(4);
+ unsorted_keys[5] = T(0); unsorted_values[5] = T(5);
+ unsorted_keys[6] = T(4); unsorted_values[6] = T(6);
+
+ sorted_keys.resize(7);
+ sorted_values.resize(7);
+ sorted_keys[0] = T(0); sorted_values[1] = T(0);
+ sorted_keys[1] = T(1); sorted_values[3] = T(1);
+ sorted_keys[2] = T(2); sorted_values[6] = T(2);
+ sorted_keys[3] = T(3); sorted_values[5] = T(3);
+ sorted_keys[4] = T(4); sorted_values[2] = T(4);
+ sorted_keys[5] = T(5); sorted_values[0] = T(5);
+ sorted_keys[6] = T(6); sorted_values[4] = T(6);
+}
+
+TYPED_TEST(SortByKeyTests, TestSortByKeySimple)
+{
+ using Vector = typename TestFixture::input_type;
+
+ Vector unsorted_keys, unsorted_values;
+ Vector sorted_keys, sorted_values;
+
+ InitializeSimpleKeyValueSortTest(unsorted_keys, unsorted_values, sorted_keys, sorted_values);
+
+ thrust::sort_by_key(unsorted_keys.begin(), unsorted_keys.end(), unsorted_values.begin());
+
+ ASSERT_EQ(unsorted_keys, sorted_keys);
+ ASSERT_EQ(unsorted_values, sorted_values);
+}
+
+TYPED_TEST(SortByKeyTests, TestSortAscendingKeyValue)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ 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 = h_keys;
+ thrust::device_vector<T> d_values = d_keys;
+
+ thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin(), thrust::less<T>());
+ thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin(), thrust::less<T>());
+
+ ASSERT_EQ(h_keys, d_keys);
+ ASSERT_EQ(h_values, d_values);
+ }
+}
+
+TEST(SortByKeyTests, TestSortDescendingKeyValue)
+{
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<int> h_keys = get_random_data<int>(size,
+ std::numeric_limits<int>::min(),
+ std::numeric_limits<int>::max());
+
+ thrust::device_vector<int> d_keys = h_keys;
+
+ thrust::host_vector<int> h_values = h_keys;
+ thrust::device_vector<int> d_values = d_keys;
+
+ thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin(), thrust::greater<int>());
+ thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin(), thrust::greater<int>());
+
+ ASSERT_EQ(h_keys, d_keys);
+ ASSERT_EQ(h_values, d_values);
+ }
+}
+
+TEST(SortByKeyTests, TestSortByKeyBool)
+{
+ const size_t size = 10027;
+
+ thrust::host_vector<bool> h_keys = get_random_data<bool>(size,
+ std::numeric_limits<bool>::min(),
+ std::numeric_limits<bool>::max());
+
+ thrust::host_vector<int> h_values = get_random_data<int>(size,
+ std::numeric_limits<int>::min(),
+ std::numeric_limits<int>::max());
+
+ thrust::device_vector<bool> d_keys = h_keys;
+ thrust::device_vector<int> d_values = h_values;
+
+ thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
+ thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
+
+ ASSERT_EQ(h_keys, d_keys);
+ ASSERT_EQ(h_values, d_values);
+}
+
+TEST(SortByKeyTests, TestSortByKeyBoolDescending)
+{
+ const size_t size = 10027;
+
+ thrust::host_vector<bool> h_keys = get_random_data<bool>(size,
+ std::numeric_limits<bool>::min(),
+ std::numeric_limits<bool>::max());
+
+ thrust::host_vector<int> h_values = get_random_data<int>(size,
+ std::numeric_limits<int>::min(),
+ std::numeric_limits<int>::max());
+
+ thrust::device_vector<bool> d_keys = h_keys;
+ thrust::device_vector<int> d_values = h_values;
+
+ thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin(), thrust::greater<bool>());
+ thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin(), thrust::greater<bool>());
+
+ ASSERT_EQ(h_keys, d_keys);
+ ASSERT_EQ(h_values, d_values);
+}

ADD
test/test_sort_by_key_variable_bits.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.
+
+// Google Test
+#include <gtest/gtest.h>
+
+#include <thrust/sort.h>
+
+#include <algorithm>
+
+#include "test_utils.hpp"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class SortByKeyVariableTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<thrust::detail::uint8_t>,
+ Params<thrust::detail::uint16_t>,
+ Params<thrust::detail::uint32_t>,
+ Params<thrust::detail::uint64_t>
+> SortByKeyVariableTestsParams;
+
+TYPED_TEST_CASE(SortByKeyVariableTests, SortByKeyVariableTestsParams);
+
+TYPED_TEST(SortByKeyVariableTests, TestSortVariableBits)
+{
+ using T = typename TestFixture::input_type;
+
+ for (auto size : get_sizes())
+ {
+ for(size_t num_bits = 0; num_bits < 8 * sizeof(T); num_bits += 3)
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<T> h_keys = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ const T mask = (1 << num_bits) - 1;
+ for(size_t i = 0; i < size; i++)
+ h_keys[i] &= mask;
+
+ thrust::host_vector<T> reference = h_keys;
+ thrust::device_vector<T> d_keys = h_keys;
+
+ thrust::host_vector<T> h_values = h_keys;
+ thrust::device_vector<T> d_values = d_keys;
+
+ std::sort(reference.begin(), reference.end());
+
+ thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
+ thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());
+
+ ASSERT_EQ(reference, h_keys);
+ ASSERT_EQ(reference, h_values);
+
+ ASSERT_EQ(h_keys, d_keys);
+ ASSERT_EQ(h_values, d_values);
+ }
+ }
+}

ADD
test/test_sort_permutation_iterator.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.
+
+// Google Test
+#include <gtest/gtest.h>
+
+#include <thrust/sort.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/transform_iterator.h>
+#include <thrust/iterator/permutation_iterator.h>
+#include <thrust/functional.h>
+
+#include "test_utils.hpp"
+
+template< class InputType >
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class SortPermutationIteratorsTests : 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::host_vector<float>>,
+ Params<thrust::host_vector<double>>,
+ 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>>,
+ Params<thrust::device_vector<float>>,
+ Params<thrust::device_vector<double>>
+> SortPermutationIteratorsTestsParams;
+
+TYPED_TEST_CASE(SortPermutationIteratorsTests, SortPermutationIteratorsTestsParams);
+
+template <typename Iterator>
+class strided_range
+{
+ public:
+
+ using difference_type = typename thrust::iterator_difference<Iterator>::type;
+
+ struct stride_functor : public thrust::unary_function<difference_type,difference_type>
+ {
+ difference_type stride;
+
+ stride_functor(difference_type stride)
+ : stride(stride) {}
+
+ __host__ __device__
+ difference_type operator()(const difference_type& i) const
+ {
+ return stride * i;
+ }
+ };
+
+ using CountingIterator = typename thrust::counting_iterator<difference_type>;
+ using TransformIterator = typename thrust::transform_iterator<stride_functor, CountingIterator>;
+ using PermutationIterator = typename thrust::permutation_iterator<Iterator,TransformIterator>;
+
+ // type of the strided_range iterator
+ using iterator = PermutationIterator;
+
+ // construct strided_range for the range [first,last)
+ strided_range(Iterator first, Iterator last, difference_type stride)
+ : first(first), last(last), stride(stride) {}
+
+ iterator begin(void) const
+ {
+ return PermutationIterator(first, TransformIterator(CountingIterator(0), stride_functor(stride)));
+ }
+
+ iterator end(void) const
+ {
+ return begin() + ((last - first) + (stride - 1)) / stride;
+ }
+
+ protected:
+ Iterator first;
+ Iterator last;
+ difference_type stride;
+};
+
+TYPED_TEST(SortPermutationIteratorsTests, TestSortPermutationIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+ using Iterator = typename Vector::iterator;
+
+ Vector A(10);
+ A[0] = T(2);
+ A[1] = T(9);
+ A[2] = T(0);
+ A[3] = T(1);
+ A[4] = T(5);
+ A[5] = T(3);
+ A[6] = T(8);
+ A[7] = T(6);
+ A[8] = T(7);
+ A[9] = T(4);
+
+ strided_range<Iterator> S(A.begin(), A.end(), 2);
+
+ thrust::sort(S.begin(), S.end());
+
+ ASSERT_EQ(A[0], T(0));
+ ASSERT_EQ(A[1], T(9));
+ ASSERT_EQ(A[2], T(2));
+ ASSERT_EQ(A[3], T(1));
+ ASSERT_EQ(A[4], T(5));
+ ASSERT_EQ(A[5], T(3));
+ ASSERT_EQ(A[6], T(7));
+ ASSERT_EQ(A[7], T(6));
+ ASSERT_EQ(A[8], T(8));
+ ASSERT_EQ(A[9], T(4));
+}
+
+TYPED_TEST(SortPermutationIteratorsTests, TestStableSortPermutationIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+ using Iterator = typename Vector::iterator;
+
+ Vector A(10);
+ A[0] = T(2);
+ A[1] = T(9);
+ A[2] = T(0);
+ A[3] = T(1);
+ A[4] = T(5);
+ A[5] = T(3);
+ A[6] = T(8);
+ A[7] = T(6);
+ A[8] = T(7);
+ A[9] = T(4);
+
+ strided_range<Iterator> S(A.begin(), A.end(), 2);
+
+ thrust::stable_sort(S.begin(), S.end());
+
+ ASSERT_EQ(A[0], T(0));
+ ASSERT_EQ(A[1], T(9));
+ ASSERT_EQ(A[2], T(2));
+ ASSERT_EQ(A[3], T(1));
+ ASSERT_EQ(A[4], T(5));
+ ASSERT_EQ(A[5], T(3));
+ ASSERT_EQ(A[6], T(7));
+ ASSERT_EQ(A[7], T(6));
+ ASSERT_EQ(A[8], T(8));
+ ASSERT_EQ(A[9], T(4));
+}
+
+TYPED_TEST(SortPermutationIteratorsTests, TestSortByKeyPermutationIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using ValueType = typename Vector::value_type;
+ using Iterator = typename Vector::iterator;
+
+ Vector A(10), B(10);
+ A[0] = ValueType(2); B[0] = ValueType(0);
+ A[1] = ValueType(9); B[1] = ValueType(1);
+ A[2] = ValueType(0); B[2] = ValueType(2);
+ A[3] = ValueType(1); B[3] = ValueType(3);
+ A[4] = ValueType(5); B[4] = ValueType(4);
+ A[5] = ValueType(3); B[5] = ValueType(5);
+ A[6] = ValueType(8); B[6] = ValueType(6);
+ A[7] = ValueType(6); B[7] = ValueType(7);
+ A[8] = ValueType(7); B[8] = ValueType(8);
+ A[9] = ValueType(4); B[9] = ValueType(9);
+
+ strided_range<Iterator> S(A.begin(), A.end(), 2);
+ strided_range<Iterator> T(B.begin(), B.end(), 2);
+
+ thrust::sort_by_key(S.begin(), S.end(), T.begin());
+
+ ASSERT_EQ(A[0], ValueType(0));
+ ASSERT_EQ(A[1], ValueType(9));
+ ASSERT_EQ(A[2], ValueType(2));
+ ASSERT_EQ(A[3], ValueType(1));
+ ASSERT_EQ(A[4], ValueType(5));
+ ASSERT_EQ(A[5], ValueType(3));
+ ASSERT_EQ(A[6], ValueType(7));
+ ASSERT_EQ(A[7], ValueType(6));
+ ASSERT_EQ(A[8], ValueType(8));
+ ASSERT_EQ(A[9], ValueType(4));
+
+ ASSERT_EQ(B[0], ValueType(2));
+ ASSERT_EQ(B[1], ValueType(1));
+ ASSERT_EQ(B[2], ValueType(0));
+ ASSERT_EQ(B[3], ValueType(3));
+ ASSERT_EQ(B[4], ValueType(4));
+ ASSERT_EQ(B[5], ValueType(5));
+ ASSERT_EQ(B[6], ValueType(8));
+ ASSERT_EQ(B[7], ValueType(7));
+ ASSERT_EQ(B[8], ValueType(6));
+ ASSERT_EQ(B[9], ValueType(9));
+}
+
+TYPED_TEST(SortPermutationIteratorsTests, TestStableSortByKeyPermutationIterator)
+{
+ using Vector = typename TestFixture::input_type;
+ using ValueType = typename Vector::value_type;
+ using Iterator = typename Vector::iterator;
+
+ Vector A(10), B(10);
+ A[0] = ValueType(2); B[0] = ValueType(0);
+ A[1] = ValueType(9); B[1] = ValueType(1);
+ A[2] = ValueType(0); B[2] = ValueType(2);
+ A[3] = ValueType(1); B[3] = ValueType(3);
+ A[4] = ValueType(5); B[4] = ValueType(4);
+ A[5] = ValueType(3); B[5] = ValueType(5);
+ A[6] = ValueType(8); B[6] = ValueType(6);
+ A[7] = ValueType(6); B[7] = ValueType(7);
+ A[8] = ValueType(7); B[8] = ValueType(8);
+ A[9] = ValueType(4); B[9] = ValueType(9);
+
+ strided_range<Iterator> S(A.begin(), A.end(), 2);
+ strided_range<Iterator> T(B.begin(), B.end(), 2);
+
+ thrust::stable_sort_by_key(S.begin(), S.end(), T.begin());
+
+ ASSERT_EQ(A[0], ValueType(0));
+ ASSERT_EQ(A[1], ValueType(9));
+ ASSERT_EQ(A[2], ValueType(2));
+ ASSERT_EQ(A[3], ValueType(1));
+ ASSERT_EQ(A[4], ValueType(5));
+ ASSERT_EQ(A[5], ValueType(3));
+ ASSERT_EQ(A[6], ValueType(7));
+ ASSERT_EQ(A[7], ValueType(6));
+ ASSERT_EQ(A[8], ValueType(8));
+ ASSERT_EQ(A[9], ValueType(4));
+
+ ASSERT_EQ(B[0], ValueType(2));
+ ASSERT_EQ(B[1], ValueType(1));
+ ASSERT_EQ(B[2], ValueType(0));
+ ASSERT_EQ(B[3], ValueType(3));
+ ASSERT_EQ(B[4], ValueType(4));
+ ASSERT_EQ(B[5], ValueType(5));
+ ASSERT_EQ(B[6], ValueType(8));
+ ASSERT_EQ(B[7], ValueType(7));
+ ASSERT_EQ(B[8], ValueType(6));
+ ASSERT_EQ(B[9], ValueType(9));
+}

ADD
test/test_sort_variables.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 <thrust/sort.h>
+
+#include <algorithm>
+
+#include "test_utils.hpp"
+
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class SortByKeyVariableTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
+typedef ::testing::Types<
+ Params<short>,
+ Params<int>,
+ Params<long long>,
+ Params<unsigned short>,
+ Params<unsigned int>,
+ Params<unsigned long long>,
+ Params<thrust::detail::uint8_t>,
+ Params<thrust::detail::uint16_t>,
+ Params<thrust::detail::uint32_t>,
+ Params<thrust::detail::uint64_t>
+> SortByKeyVariableTestsParams;
+
+TYPED_TEST_CASE(SortByKeyVariableTests, SortByKeyVariableTestsParams);
+
+TYPED_TEST(SortByKeyVariableTests, TestSortVariableBits)
+{
+ using T = typename TestFixture::input_type;
+
+ for (auto size : get_sizes())
+ {
+ for(size_t num_bits = 0; num_bits < 8 * sizeof(T); num_bits += 3)
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<T> h_keys = get_random_data<T>(size,
+ std::numeric_limits<T>::min(),
+ std::numeric_limits<T>::max());
+
+ size_t mask = (1 << num_bits) - 1;
+ for(size_t i = 0; i < size; i++)
+ h_keys[i] &= mask;
+
+ thrust::host_vector<T> reference = h_keys;
+ thrust::device_vector<T> d_keys = h_keys;
+
+ std::sort(reference.begin(), reference.end());
+
+ thrust::sort(h_keys.begin(), h_keys.end());
+ thrust::sort(d_keys.begin(), d_keys.end());
+
+ ASSERT_EQ(reference, h_keys);
+ ASSERT_EQ(h_keys, d_keys);
+ }
+ }
+}

MODIFY
test/test_utils.hpp
test/test_utils.hpp
@@ -38,9 +38,21 @@ std::vector<size_t> get_sizes()
return sizes;
}
+template<class T>
+inline auto get_random_data(size_t size, T, T)
+ -> typename std::enable_if<std::is_same<T,bool>::value, thrust::host_vector<T>>::type
+{
+ std::random_device rd;
+ std::default_random_engine gen(rd());
+ std::bernoulli_distribution distribution(0.5);
+ 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_integral<T>::value, thrust::host_vector<T>>::type
+ -> typename std::enable_if<rocprim::is_integral<T>::value && !std::is_same<T,bool>::value, thrust::host_vector<T>>::type
{
std::random_device rd;
std::default_random_engine gen(rd());

Implemented sort tests