Extended unique tests
MODIFY
test/CMakeLists.txt
test/CMakeLists.txt
@@ -106,6 +106,7 @@ add_hipthrust_test("thrust.hip.tuple_sort" test_tuple_sort.cpp)
add_hipthrust_test("thrust.hip.tuple_transform" test_tuple_transform.cpp)
add_hipthrust_test("thrust.hip.uninitialized_copy" test_uninitialized_copy.cpp)
add_hipthrust_test("thrust.hip.uninitialized_fill" test_uninitialized_fill.cpp)
+add_hipthrust_test("thrust.hip.unique" test_unique.cpp)
add_hipthrust_test("thrust.hip.unique_by_key" test_unique_by_key.cpp)
add_hipthrust_test("thrust.hip.vector" test_vector.cpp)
add_hipthrust_test("thrust.hip.vector_insert" test_vector_insert.cpp)

ADD
test/test_unique.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>
+
+#include <thrust/unique.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/discard_iterator.h>
+#include <thrust/iterator/retag.h>
+
+// HIP API
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <hip/hip_runtime_api.h>
+
+#define HIP_CHECK(condition) ASSERT_EQ(condition, hipSuccess)
+#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+#include "test_utils.hpp"
+#include "test_assertions.hpp"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class UniqueTests : 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>>
+> UniqueTestsParams;
+
+TYPED_TEST_CASE(UniqueTests, UniqueTestsParams);
+
+template<class Params>
+class UniqueIntegralTests : 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>
+> UniqueIntegralTestsParams;
+
+TYPED_TEST_CASE(UniqueIntegralTests, UniqueIntegralTestsParams);
+
+template <typename ForwardIterator>
+ForwardIterator unique(my_system &system,
+ ForwardIterator first,
+ ForwardIterator)
+{
+ system.validate_dispatch();
+ return first;
+}
+
+TEST(UniqueTests, TestUniqueDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::unique(sys, vec.begin(), vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template <typename ForwardIterator>
+ForwardIterator unique(my_tag,
+ ForwardIterator first,
+ ForwardIterator)
+{
+ *first = 13;
+ return first;
+}
+
+TEST(UniqueTests, TestUniqueDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::unique(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()));
+
+ ASSERT_EQ(13, vec.front());
+}
+
+template <typename InputIterator,
+ typename OutputIterator>
+OutputIterator unique_copy(my_system &system,
+ InputIterator,
+ InputIterator,
+ OutputIterator result)
+{
+ system.validate_dispatch();
+ return result;
+}
+
+TEST(UniqueTests, TestUniqueCopyDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::unique_copy(sys, vec.begin(), vec.begin(), vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template <typename InputIterator,
+ typename OutputIterator>
+OutputIterator unique_copy(my_tag,
+ InputIterator,
+ InputIterator,
+ OutputIterator result)
+{
+ *result = 13;
+ return result;
+}
+
+TEST(UniqueTests, TestUniqueCopyDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::unique_copy(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 is_equal_div_10_unique
+{
+ __host__ __device__
+ bool operator()(const T x, const T& y) const { return ((int) x / 10) == ((int) y / 10); }
+};
+
+TYPED_TEST(UniqueTests, TestUniqueSimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(10);
+ data[0] = 11;
+ data[1] = 11;
+ data[2] = 12;
+ data[3] = 20;
+ data[4] = 29;
+ data[5] = 21;
+ data[6] = 21;
+ data[7] = 31;
+ data[8] = 31;
+ data[9] = 37;
+
+ typename Vector::iterator new_last;
+
+ new_last = thrust::unique(data.begin(), data.end());
+
+ ASSERT_EQ(new_last - data.begin(), 7);
+ ASSERT_EQ(data[0], 11);
+ ASSERT_EQ(data[1], 12);
+ ASSERT_EQ(data[2], 20);
+ ASSERT_EQ(data[3], 29);
+ ASSERT_EQ(data[4], 21);
+ ASSERT_EQ(data[5], 31);
+ ASSERT_EQ(data[6], 37);
+
+ new_last = thrust::unique(data.begin(), new_last, is_equal_div_10_unique<T>());
+
+ ASSERT_EQ(new_last - data.begin(), 3);
+ ASSERT_EQ(data[0], 11);
+ ASSERT_EQ(data[1], 20);
+ ASSERT_EQ(data[2], 31);
+}
+
+TYPED_TEST(UniqueIntegralTests, TestUnique)
+{
+ using T = typename TestFixture::input_type;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ 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;
+
+ typename thrust::host_vector<T>::iterator h_new_last;
+ typename thrust::device_vector<T>::iterator d_new_last;
+
+ h_new_last = thrust::unique(h_data.begin(), h_data.end());
+ d_new_last = thrust::unique(d_data.begin(), d_data.end());
+
+ ASSERT_EQ(h_new_last - h_data.begin(), d_new_last - d_data.begin());
+
+ h_data.resize(h_new_last - h_data.begin());
+ d_data.resize(d_new_last - d_data.begin());
+
+ ASSERT_EQ(h_data, d_data);
+ }
+}
+
+TYPED_TEST(UniqueTests, TestUniqueCopySimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector data(10);
+ data[0] = 11;
+ data[1] = 11;
+ data[2] = 12;
+ data[3] = 20;
+ data[4] = 29;
+ data[5] = 21;
+ data[6] = 21;
+ data[7] = 31;
+ data[8] = 31;
+ data[9] = 37;
+
+ Vector output(10, -1);
+
+ typename Vector::iterator new_last;
+
+ new_last = thrust::unique_copy(data.begin(), data.end(), output.begin());
+
+ ASSERT_EQ(new_last - output.begin(), 7);
+ ASSERT_EQ(output[0], 11);
+ ASSERT_EQ(output[1], 12);
+ ASSERT_EQ(output[2], 20);
+ ASSERT_EQ(output[3], 29);
+ ASSERT_EQ(output[4], 21);
+ ASSERT_EQ(output[5], 31);
+ ASSERT_EQ(output[6], 37);
+
+ new_last = thrust::unique_copy(output.begin(), new_last, data.begin(), is_equal_div_10_unique<T>());
+
+ ASSERT_EQ(new_last - data.begin(), 3);
+ ASSERT_EQ(data[0], 11);
+ ASSERT_EQ(data[1], 20);
+ ASSERT_EQ(data[2], 31);
+}
+
+TYPED_TEST(UniqueIntegralTests, TestUniqueCopy)
+{
+ using T = typename TestFixture::input_type;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ 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::host_vector<T> h_output(size);
+ thrust::device_vector<T> d_output(size);
+
+ typename thrust::host_vector<T>::iterator h_new_last;
+ typename thrust::device_vector<T>::iterator d_new_last;
+
+ h_new_last = thrust::unique_copy(h_data.begin(), h_data.end(), h_output.begin());
+ d_new_last = thrust::unique_copy(d_data.begin(), d_data.end(), d_output.begin());
+
+ ASSERT_EQ(h_new_last - h_output.begin(), d_new_last - d_output.begin());
+
+ h_data.resize(h_new_last - h_output.begin());
+ d_data.resize(d_new_last - d_output.begin());
+
+ ASSERT_EQ(h_output, d_output);
+ }
+}
+
+TYPED_TEST(UniqueIntegralTests, TestUniqueCopyToDiscardIterator)
+{
+ using T = typename TestFixture::input_type;
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ 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::host_vector<T> h_unique = h_data;
+ h_unique.erase(thrust::unique(h_unique.begin(), h_unique.end()), h_unique.end());
+
+ thrust::discard_iterator<> reference(h_unique.size());
+
+ typename thrust::host_vector<T>::iterator h_new_last;
+ typename thrust::device_vector<T>::iterator d_new_last;
+
+ thrust::discard_iterator<> h_result =
+ thrust::unique_copy(h_data.begin(), h_data.end(), thrust::make_discard_iterator());
+
+ thrust::discard_iterator<> d_result =
+ thrust::unique_copy(d_data.begin(), d_data.end(), thrust::make_discard_iterator());
+
+ ASSERT_EQ_QUIET(reference, h_result);
+ ASSERT_EQ_QUIET(reference, d_result);
+ }
+}

MODIFY
test/test_unique_by_key.cpp
test/test_unique_by_key.cpp
#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include "test_utils.hpp"
+#include "test_assertions.hpp"
template<
class InputType
@@ -73,6 +74,23 @@ typedef ::testing::Types<
TYPED_TEST_CASE(UniqueByKeyTests, UniqueByKeyTestsParams);
+template<class Params>
+class UniqueByKeyIntegralTests : 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>
+> UniqueByKeyIntegralTestsParams;
+
+TYPED_TEST_CASE(UniqueByKeyIntegralTests, UniqueByKeyIntegralTestsParams);
template <typename ForwardIterator1,
typename ForwardIterator2>
@@ -309,3 +327,207 @@ TYPED_TEST(UniqueByKeyTests, TestUniqueCopyByKeySimple)
ASSERT_EQ(output_values[1], 2);
ASSERT_EQ(output_values[2], 7);
}
+
+TYPED_TEST(UniqueByKeyIntegralTests, TestUniqueByKey)
+{
+ using K = typename TestFixture::input_type;
+ using V = unsigned int; // ValueType
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+
+
+ thrust::host_vector<K> h_keys = get_random_data<K>(size,
+ std::numeric_limits<K>::min(),
+ std::numeric_limits<K>::max());
+
+ thrust::host_vector<V> h_vals = get_random_data<V>(size,
+ std::numeric_limits<V>::min(),
+ std::numeric_limits<V>::max());;
+ thrust::device_vector<K> d_keys = h_keys;
+ thrust::device_vector<V> d_vals = h_vals;
+
+ using HostKeyIterator = typename thrust::host_vector<K>::iterator;
+ using HostValIterator = typename thrust::host_vector<V>::iterator;
+ using DeviceKeyIterator = typename thrust::device_vector<K>::iterator;
+ using DeviceValIterator = typename thrust::device_vector<V>::iterator;
+
+ using HostIteratorPair = typename thrust::pair<HostKeyIterator, HostValIterator>;
+ using DeviceIteratorPair = typename thrust::pair<DeviceKeyIterator,DeviceValIterator>;
+
+ HostIteratorPair h_last = thrust::unique_by_key(h_keys.begin(), h_keys.end(), h_vals.begin());
+ DeviceIteratorPair d_last = thrust::unique_by_key(d_keys.begin(), d_keys.end(), d_vals.begin());
+
+ ASSERT_EQ(h_last.first - h_keys.begin(), d_last.first - d_keys.begin());
+ ASSERT_EQ(h_last.second - h_vals.begin(), d_last.second - d_vals.begin());
+
+ size_t N = h_last.first - h_keys.begin();
+
+ h_keys.resize(N);
+ h_vals.resize(N);
+ d_keys.resize(N);
+ d_vals.resize(N);
+
+ ASSERT_EQ(h_keys, d_keys);
+ ASSERT_EQ(h_vals, d_vals);
+ }
+}
+
+TYPED_TEST(UniqueByKeyIntegralTests, TestUniqueCopyByKey)
+{
+ using K = typename TestFixture::input_type;
+ using V = unsigned int; // ValueType
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+
+
+ thrust::host_vector<K> h_keys = get_random_data<K>(size,
+ std::numeric_limits<K>::min(),
+ std::numeric_limits<K>::max());
+
+ thrust::host_vector<V> h_vals = get_random_data<V>(size,
+ std::numeric_limits<V>::min(),
+ std::numeric_limits<V>::max());;
+ thrust::device_vector<K> d_keys = h_keys;
+ thrust::device_vector<V> d_vals = h_vals;
+
+ thrust::host_vector<K> h_keys_output(size);
+ thrust::host_vector<V> h_vals_output(size);
+ thrust::device_vector<K> d_keys_output(size);
+ thrust::device_vector<V> d_vals_output(size);
+
+ using HostKeyIterator = typename thrust::host_vector<K>::iterator;
+ using HostValIterator = typename thrust::host_vector<V>::iterator;
+ using DeviceKeyIterator = typename thrust::device_vector<K>::iterator;
+ using DeviceValIterator = typename thrust::device_vector<V>::iterator;
+
+ using HostIteratorPair = typename thrust::pair<HostKeyIterator, HostValIterator>;
+ using DeviceIteratorPair = typename thrust::pair<DeviceKeyIterator,DeviceValIterator>;
+
+ HostIteratorPair h_last = thrust::unique_by_key_copy(h_keys.begin(), h_keys.end(), h_vals.begin(), h_keys_output.begin(), h_vals_output.begin());
+ DeviceIteratorPair d_last = thrust::unique_by_key_copy(d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys_output.begin(), d_vals_output.begin());
+
+ ASSERT_EQ(h_last.first - h_keys_output.begin(), d_last.first - d_keys_output.begin());
+ ASSERT_EQ(h_last.second - h_vals_output.begin(), d_last.second - d_vals_output.begin());
+
+ size_t N = h_last.first - h_keys_output.begin();
+
+ h_keys_output.resize(N);
+ h_vals_output.resize(N);
+ d_keys_output.resize(N);
+ d_vals_output.resize(N);
+
+ ASSERT_EQ(h_keys_output, d_keys_output);
+ ASSERT_EQ(h_vals_output, d_vals_output);
+ }
+}
+
+TYPED_TEST(UniqueByKeyIntegralTests, TestUniqueCopyByKeyToDiscardIterator)
+{
+ using K = typename TestFixture::input_type;
+ using V = unsigned int; // ValueType
+
+ for (auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+
+ thrust::host_vector<K> h_keys = get_random_data<K>(size,
+ std::numeric_limits<K>::min(),
+ std::numeric_limits<K>::max());
+
+ thrust::host_vector<V> h_vals = get_random_data<V>(size,
+ std::numeric_limits<V>::min(),
+ std::numeric_limits<V>::max());;
+ thrust::device_vector<K> d_keys = h_keys;
+ thrust::device_vector<V> d_vals = h_vals;
+
+ thrust::host_vector<V> h_vals_output(size);
+ thrust::device_vector<V> d_vals_output(size);
+
+ thrust::host_vector<K> h_keys_output(size);
+ thrust::device_vector<K> d_keys_output(size);
+
+ thrust::host_vector<K> h_unique_keys = h_keys;
+ h_unique_keys.erase(thrust::unique(h_unique_keys.begin(), h_unique_keys.end()), h_unique_keys.end());
+
+ size_t num_unique_keys = h_unique_keys.size();
+
+
+ // mask both outputs
+ thrust::pair<thrust::discard_iterator<>, thrust::discard_iterator<> > h_result1 =
+ thrust::unique_by_key_copy(h_keys.begin(), h_keys.end(),
+ h_vals.begin(),
+ thrust::make_discard_iterator(),
+ thrust::make_discard_iterator());
+
+ thrust::pair<thrust::discard_iterator<>, thrust::discard_iterator<> > d_result1 =
+ thrust::unique_by_key_copy(d_keys.begin(), d_keys.end(),
+ d_vals.begin(),
+ thrust::make_discard_iterator(),
+ thrust::make_discard_iterator());
+
+ thrust::pair<thrust::discard_iterator<>, thrust::discard_iterator<> > reference1 =
+ thrust::make_pair(thrust::make_discard_iterator(num_unique_keys),
+ thrust::make_discard_iterator(num_unique_keys));
+
+ ASSERT_EQ_QUIET(reference1, h_result1);
+ ASSERT_EQ_QUIET(reference1, d_result1);
+
+
+ // mask values output
+ thrust::pair<typename thrust::host_vector<K>::iterator, thrust::discard_iterator<> > h_result2 =
+ thrust::unique_by_key_copy(h_keys.begin(), h_keys.end(),
+ h_vals.begin(),
+ h_keys_output.begin(),
+ thrust::make_discard_iterator());
+
+ thrust::pair<typename thrust::device_vector<K>::iterator, thrust::discard_iterator<> > d_result2 =
+ thrust::unique_by_key_copy(d_keys.begin(), d_keys.end(),
+ d_vals.begin(),
+ d_keys_output.begin(),
+ thrust::make_discard_iterator());
+
+ thrust::pair<typename thrust::host_vector<K>::iterator, thrust::discard_iterator<> > h_reference2 =
+ thrust::make_pair(h_keys_output.begin() + num_unique_keys,
+ thrust::make_discard_iterator(num_unique_keys));
+
+ thrust::pair<typename thrust::device_vector<K>::iterator, thrust::discard_iterator<> > d_reference2 =
+ thrust::make_pair(d_keys_output.begin() + num_unique_keys,
+ thrust::make_discard_iterator(num_unique_keys));
+
+ ASSERT_EQ(h_keys_output, d_keys_output);
+ ASSERT_EQ_QUIET(h_reference2, h_result2);
+ ASSERT_EQ_QUIET(d_reference2, d_result2);
+
+
+ // mask keys output
+ thrust::pair<thrust::discard_iterator<>, typename thrust::host_vector<V>::iterator> h_result3 =
+ thrust::unique_by_key_copy(h_keys.begin(), h_keys.end(),
+ h_vals.begin(),
+ thrust::make_discard_iterator(),
+ h_vals_output.begin());
+
+ thrust::pair<thrust::discard_iterator<>, typename thrust::device_vector<V>::iterator> d_result3 =
+ thrust::unique_by_key_copy(d_keys.begin(), d_keys.end(),
+ d_vals.begin(),
+ thrust::make_discard_iterator(),
+ d_vals_output.begin());
+
+ thrust::pair<thrust::discard_iterator<>, typename thrust::host_vector<V>::iterator> h_reference3 =
+ thrust::make_pair(thrust::make_discard_iterator(num_unique_keys),
+ h_vals_output.begin() + num_unique_keys);
+
+ thrust::pair<thrust::discard_iterator<>, typename thrust::device_vector<V>::iterator> d_reference3 =
+ thrust::make_pair(thrust::make_discard_iterator(num_unique_keys),
+ d_vals_output.begin() + num_unique_keys);
+
+ ASSERT_EQ(h_vals_output, d_vals_output);
+ ASSERT_EQ_QUIET(h_reference3, h_result3);
+ ASSERT_EQ_QUIET(d_reference3, d_result3);
+ }
+}

RENAME
thrust/system/hip/detail/unqiue.h
thrust/system/hip/detail/unique.h
Extended unique tests