Add unique_by_key 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_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_hipthrust_test("thrust.hip.vector_manipulation" test_vector_manipulation.cpp)

ADD
test/test_unique_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>
+
+#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"
+
+template<
+ class InputType
+>
+struct Params
+{
+ using input_type = InputType;
+};
+
+template<class Params>
+class UniqueByKeyTests : 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>>
+> UniqueByKeyTestsParams;
+
+TYPED_TEST_CASE(UniqueByKeyTests, UniqueByKeyTestsParams);
+
+
+template <typename ForwardIterator1,
+ typename ForwardIterator2>
+thrust::pair<ForwardIterator1,ForwardIterator2>
+unique_by_key(my_system &system,
+ ForwardIterator1 keys_first,
+ ForwardIterator1,
+ ForwardIterator2 values_first)
+{
+ system.validate_dispatch();
+ return thrust::make_pair(keys_first,values_first);
+}
+
+TEST(UniqueByKeyTests, TestUniqueByKeyDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::unique_by_key(sys, vec.begin(), vec.begin(), vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template <typename ForwardIterator1,
+ typename ForwardIterator2>
+thrust::pair<ForwardIterator1,ForwardIterator2>
+unique_by_key(my_tag,
+ ForwardIterator1 keys_first,
+ ForwardIterator1,
+ ForwardIterator2 values_first)
+{
+ *keys_first = 13;
+ return thrust::make_pair(keys_first,values_first);
+}
+
+TEST(UniqueByKeyTests, TestUniqueByKeyDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::unique_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 InputIterator1,
+ typename InputIterator2,
+ typename OutputIterator1,
+ typename OutputIterator2>
+thrust::pair<OutputIterator1,OutputIterator2>
+unique_by_key_copy(my_system &system,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ OutputIterator1 keys_output,
+ OutputIterator2 values_output)
+{
+ system.validate_dispatch();
+ return thrust::make_pair(keys_output, values_output);
+}
+
+TEST(UniqueByKeyTests, TestUniqueByKeyCopyDispatchExplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ my_system sys(0);
+ thrust::unique_by_key_copy(sys,
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ vec.begin(),
+ vec.begin());
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template <typename InputIterator1,
+ typename InputIterator2,
+ typename OutputIterator1,
+ typename OutputIterator2>
+thrust::pair<OutputIterator1,OutputIterator2>
+unique_by_key_copy(my_tag,
+ InputIterator1,
+ InputIterator1,
+ InputIterator2,
+ OutputIterator1 keys_output,
+ OutputIterator2 values_output)
+{
+ *keys_output = 13;
+ return thrust::make_pair(keys_output, values_output);
+}
+
+TEST(UniqueByKeyTests, TestUniqueByKeyCopyDispatchImplicit)
+{
+ thrust::device_vector<int> vec(1);
+
+ thrust::unique_by_key_copy(thrust::retag<my_tag>(vec.begin()),
+ thrust::retag<my_tag>(vec.begin()),
+ 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); }
+};
+
+template <typename Vector>
+void initialize_keys(Vector& keys)
+{
+ keys.resize(9);
+ keys[0] = 11;
+ keys[1] = 11;
+ keys[2] = 21;
+ keys[3] = 20;
+ keys[4] = 21;
+ keys[5] = 21;
+ keys[6] = 21;
+ keys[7] = 37;
+ keys[8] = 37;
+}
+
+template <typename Vector>
+void initialize_values(Vector& values)
+{
+ values.resize(9);
+ values[0] = 0;
+ values[1] = 1;
+ values[2] = 2;
+ values[3] = 3;
+ values[4] = 4;
+ values[5] = 5;
+ values[6] = 6;
+ values[7] = 7;
+ values[8] = 8;
+}
+
+
+TYPED_TEST(UniqueByKeyTests, TestUniqueByKeySimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector keys;
+ Vector values;
+
+ typename thrust::pair<typename Vector::iterator, typename Vector::iterator> new_last;
+
+ // basic test
+ initialize_keys(keys); initialize_values(values);
+
+ new_last = thrust::unique_by_key(keys.begin(), keys.end(), values.begin());
+
+ ASSERT_EQ(new_last.first - keys.begin(), 5);
+ ASSERT_EQ(new_last.second - values.begin(), 5);
+ ASSERT_EQ(keys[0], 11);
+ ASSERT_EQ(keys[1], 21);
+ ASSERT_EQ(keys[2], 20);
+ ASSERT_EQ(keys[3], 21);
+ ASSERT_EQ(keys[4], 37);
+
+ ASSERT_EQ(values[0], 0);
+ ASSERT_EQ(values[1], 2);
+ ASSERT_EQ(values[2], 3);
+ ASSERT_EQ(values[3], 4);
+ ASSERT_EQ(values[4], 7);
+
+ // test BinaryPredicate
+ initialize_keys(keys); initialize_values(values);
+
+ new_last = thrust::unique_by_key(keys.begin(), keys.end(), values.begin(), is_equal_div_10_unique<T>());
+
+ ASSERT_EQ(new_last.first - keys.begin(), 3);
+ ASSERT_EQ(new_last.second - values.begin(), 3);
+ ASSERT_EQ(keys[0], 11);
+ ASSERT_EQ(keys[1], 21);
+ ASSERT_EQ(keys[2], 37);
+
+ ASSERT_EQ(values[0], 0);
+ ASSERT_EQ(values[1], 2);
+ ASSERT_EQ(values[2], 7);
+}
+
+TYPED_TEST(UniqueByKeyTests, TestUniqueCopyByKeySimple)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector keys;
+ Vector values;
+
+ typename thrust::pair<typename Vector::iterator, typename Vector::iterator> new_last;
+
+ // basic test
+ initialize_keys(keys); initialize_values(values);
+
+ Vector output_keys(keys.size());
+ Vector output_values(values.size());
+
+ new_last = thrust::unique_by_key_copy(keys.begin(), keys.end(), values.begin(), output_keys.begin(), output_values.begin());
+
+ ASSERT_EQ(new_last.first - output_keys.begin(), 5);
+ ASSERT_EQ(new_last.second - output_values.begin(), 5);
+ ASSERT_EQ(output_keys[0], 11);
+ ASSERT_EQ(output_keys[1], 21);
+ ASSERT_EQ(output_keys[2], 20);
+ ASSERT_EQ(output_keys[3], 21);
+ ASSERT_EQ(output_keys[4], 37);
+
+ ASSERT_EQ(output_values[0], 0);
+ ASSERT_EQ(output_values[1], 2);
+ ASSERT_EQ(output_values[2], 3);
+ ASSERT_EQ(output_values[3], 4);
+ ASSERT_EQ(output_values[4], 7);
+
+ // test BinaryPredicate
+ initialize_keys(keys); initialize_values(values);
+
+ new_last = thrust::unique_by_key_copy(keys.begin(), keys.end(), values.begin(), output_keys.begin(), output_values.begin(), is_equal_div_10_unique<T>());
+
+ ASSERT_EQ(new_last.first - output_keys.begin(), 3);
+ ASSERT_EQ(new_last.second - output_values.begin(), 3);
+ ASSERT_EQ(output_keys[0], 11);
+ ASSERT_EQ(output_keys[1], 21);
+ ASSERT_EQ(output_keys[2], 37);
+
+ ASSERT_EQ(output_values[0], 0);
+ ASSERT_EQ(output_values[1], 2);
+ ASSERT_EQ(output_values[2], 7);
+}

ADD
thrust/system/hip/detail/unique_by_key.h
+/*
+ * Copyright 2008-2013 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// this system has no special version of this algorithm
+
+/*
+
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <thrust/system/cuda/config.h>
+
+#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/get_value.h>
+#include <thrust/system/hip/detail/par_to_seq.h>
+#include <thrust/system/hip/detail/memory_buffer.h>
+#include <thrust/functional.h>
+#include <thrust/pair.h>
+#include <thrust/detail/mpl/math.h>
+#include <thrust/detail/minmax.h>
+#include <thrust/distance.h>
+
+BEGIN_NS_THRUST
+
+template <typename DerivedPolicy,
+ typename ForwardIterator1,
+ typename ForwardIterator2>
+__host__ __device__ thrust::pair<ForwardIterator1, ForwardIterator2>
+unique_by_key(
+ const thrust::detail::execution_policy_base<DerivedPolicy> &exec,
+ ForwardIterator1 keys_first,
+ ForwardIterator1 keys_last,
+ ForwardIterator2 values_first);
+template <typename DerivedPolicy,
+ typename InputIterator1,
+ typename InputIterator2,
+ typename OutputIterator1,
+ typename OutputIterator2>
+__host__ __device__ thrust::pair<OutputIterator1, OutputIterator2>
+unique_by_key_copy(
+ const thrust::detail::execution_policy_base<DerivedPolicy> &exec,
+ InputIterator1 keys_first,
+ InputIterator1 keys_last,
+ InputIterator2 values_first,
+ OutputIterator1 keys_result,
+ OutputIterator2 values_result);
+
+namespace hip_rocprim {
+
+// XXX it should be possible to unify unique & unique_by_key into a single
+// agent with various specializations, similar to what is done
+// with partition
+namespace __unique_by_key {
+
+
+ template <class Policy,
+ class KeyInputIt,
+ class ValInputIt,
+ class KeyOutputIt,
+ class ValOutputIt,
+ class BinaryPred>
+ pair<KeyOutputIt, ValOutputIt> THRUST_RUNTIME_FUNCTION
+ unique_by_key(Policy & policy,
+ KeyInputIt keys_first,
+ KeyInputIt keys_last,
+ ValInputIt values_first,
+ KeyOutputIt keys_result,
+ ValOutputIt values_result,
+ BinaryPred binary_pred)
+ {
+
+ // typedef typename iterator_traits<KeyInputIt>::difference_type size_type;
+ typedef int size_type;
+
+ size_type num_items = static_cast<size_type>(thrust::distance(keys_first, keys_last));
+ char * d_temp_storage = NULL;
+ size_t temp_storage_bytes = 0;
+ cudaStream_t stream = hip_rocprim::stream(policy);
+ size_type * d_num_selected_out = NULL;
+ bool debug_sync = THRUST_DEBUG_SYNC_FLAG;
+
+ cudaError_t status;
+ status = __unique_by_key::doit_step(d_temp_storage,
+ temp_storage_bytes,
+ keys_first,
+ values_first,
+ keys_result,
+ values_result,
+ binary_pred,
+ d_num_selected_out,
+ num_items,
+ stream,
+ debug_sync);
+ hip_rocprim::throw_on_error(status, "unique_by_key: failed on 1st step");
+
+ size_t allocation_sizes[2] = {sizeof(size_type), temp_storage_bytes};
+ void * allocations[2] = {NULL, NULL};
+
+ size_t storage_size = 0;
+ status = core::alias_storage(NULL,
+ storage_size,
+ allocations,
+ allocation_sizes);
+
+ void *ptr = hip_rocprim::get_memory_buffer(policy, storage_size);
+ hip_rocprim::throw_on_error(cudaGetLastError(),
+ "unique_by_key: failed to get memory buffer");
+
+ status = core::alias_storage(ptr,
+ storage_size,
+ allocations,
+ allocation_sizes);
+
+ d_num_selected_out = (size_type *)allocations[0];
+ d_temp_storage = (char *)allocations[1];
+
+ status = __unique_by_key::doit_step(d_temp_storage,
+ temp_storage_bytes,
+ keys_first,
+ values_first,
+ keys_result,
+ values_result,
+ binary_pred,
+ d_num_selected_out,
+ num_items,
+ stream,
+ debug_sync);
+ hip_rocprim::throw_on_error(status, "unique_by_key: failed on 2nd step");
+
+
+ status = hip_rocprim::synchronize(policy);
+ hip_rocprim::throw_on_error(status, "unique_by_key: failed to synchronize");
+
+ size_type num_selected = get_value(policy, d_num_selected_out);
+
+ hip_rocprim::return_memory_buffer(policy, ptr);
+ hip_rocprim::throw_on_error(cudaGetLastError(),
+ "unique_by_key: failed to return memory buffer");
+
+ return thrust::make_pair(keys_result + num_selected, values_result + num_selected);
+ }
+} // namespace __unique_by_key
+
+//-------------------------
+// Thrust API entry points
+//-------------------------
+
+
+__thrust_exec_check_disable__
+template <class Derived,
+ class KeyInputIt,
+ class ValInputIt,
+ class KeyOutputIt,
+ class ValOutputIt,
+ class BinaryPred>
+pair<KeyOutputIt, ValOutputIt> __host__ __device__
+unique_by_key_copy(execution_policy<Derived> &policy,
+ KeyInputIt keys_first,
+ KeyInputIt keys_last,
+ ValInputIt values_first,
+ KeyOutputIt keys_result,
+ ValOutputIt values_result,
+ BinaryPred binary_pred)
+{
+ pair<KeyOutputIt, ValOutputIt> ret = thrust::make_pair(keys_result, values_result);
+ if (__THRUST_HAS_CUDART__)
+ {
+ ret = __unique_by_key::unique_by_key(policy,
+ keys_first,
+ keys_last,
+ values_first,
+ keys_result,
+ values_result,
+ binary_pred);
+ }
+ else
+ {
+#if !__THRUST_HAS_CUDART__
+ ret = thrust::unique_by_key_copy(cvt_to_seq(derived_cast(policy)),
+ keys_first,
+ keys_last,
+ values_first,
+ keys_result,
+ values_result,
+ binary_pred);
+#endif
+ }
+ return ret;
+}
+
+template <class Derived,
+ class KeyInputIt,
+ class ValInputIt,
+ class KeyOutputIt,
+ class ValOutputIt>
+pair<KeyOutputIt, ValOutputIt> __host__ __device__
+unique_by_key_copy(execution_policy<Derived> &policy,
+ KeyInputIt keys_first,
+ KeyInputIt keys_last,
+ ValInputIt values_first,
+ KeyOutputIt keys_result,
+ ValOutputIt values_result)
+{
+ typedef typename iterator_traits<KeyInputIt>::value_type key_type;
+ return hip_rocprim::unique_by_key_copy(policy,
+ keys_first,
+ keys_last,
+ values_first,
+ keys_result,
+ values_result,
+ equal_to<key_type>());
+}
+
+template <class Derived,
+ class KeyInputIt,
+ class ValInputIt,
+ class BinaryPred>
+pair<KeyInputIt, ValInputIt> __host__ __device__
+unique_by_key(execution_policy<Derived> &policy,
+ KeyInputIt keys_first,
+ KeyInputIt keys_last,
+ ValInputIt values_first,
+ BinaryPred binary_pred)
+{
+ pair<KeyInputIt, ValInputIt> ret = thrust::make_pair(keys_first, values_first);
+ if (__THRUST_HAS_CUDART__)
+ {
+ ret = hip_rocprim::unique_by_key_copy(policy,
+ keys_first,
+ keys_last,
+ values_first,
+ keys_first,
+ values_first,
+ binary_pred);
+ }
+ else
+ {
+#if !__THRUST_HAS_CUDART__
+ ret = thrust::unique_by_key(cvt_to_seq(derived_cast(policy)),
+ keys_first,
+ keys_last,
+ values_first,
+ binary_pred);
+#endif
+ }
+ return ret;
+}
+
+template <class Derived,
+ class KeyInputIt,
+ class ValInputIt>
+pair<KeyInputIt, ValInputIt> __host__ __device__
+unique_by_key(execution_policy<Derived> &policy,
+ KeyInputIt keys_first,
+ KeyInputIt keys_last,
+ ValInputIt values_first)
+{
+ typedef typename iterator_traits<KeyInputIt>::value_type key_type;
+ return hip_rocprim::unique_by_key(policy,
+ keys_first,
+ keys_last,
+ values_first,
+ equal_to<key_type>());
+}
+
+
+
+} // namespace hip_rocprim
+END_NS_THRUST
+
+#include <thrust/memory.h>
+#include <thrust/unique.h>
+
+#endif*/

ADD
thrust/system/hip/detail/unqiue.h
+/*
+ * Copyright 2008-2013 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// this system has no special version of this algorithm

Add unique_by_key tests