From 8aecfe5e06962ec8fe32d307b821653523959190 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Sun, 13 Feb 2022 14:05:02 +0100 Subject: [PATCH 01/10] add unique_count algorithm Add a counting equivalent to unique_* algorithms that can be used to allocate the correct amount of data before actually filling it. Addresses issue #1612 --- testing/unique.cu | 89 ++++++++++++ thrust/count.h | 2 +- thrust/detail/count.h | 60 ++++++++ thrust/detail/unique.inl | 61 ++++++++ thrust/system/cuda/detail/unique.h | 40 ++++++ thrust/system/detail/generic/unique.h | 20 +++ thrust/system/detail/generic/unique.inl | 32 +++++ thrust/system/detail/sequential/unique.h | 34 +++++ thrust/system/omp/detail/unique.h | 10 ++ thrust/system/omp/detail/unique.inl | 15 ++ thrust/system/tbb/detail/unique.h | 10 ++ thrust/system/tbb/detail/unique.inl | 15 ++ thrust/unique.h | 175 +++++++++++++++++++++++ 13 files changed, 562 insertions(+), 1 deletion(-) create mode 100644 thrust/detail/count.h diff --git a/testing/unique.cu b/testing/unique.cu index 8073832df..7df2def87 100644 --- a/testing/unique.cu +++ b/testing/unique.cu @@ -95,6 +95,50 @@ void TestUniqueCopyDispatchImplicit() DECLARE_UNITTEST(TestUniqueCopyDispatchImplicit); +template +typename thrust::iterator_traits::difference_type + unique_count(my_system &system, + ForwardIterator, + ForwardIterator) +{ + system.validate_dispatch(); + return 0; +} + +void TestUniqueCountDispatchExplicit() +{ + thrust::device_vector vec(1); + + my_system sys(0); + thrust::unique_count(sys, vec.begin(), vec.begin()); + + ASSERT_EQUAL(true, sys.is_valid()); +} +DECLARE_UNITTEST(TestUniqueCountDispatchExplicit); + + +template +typename thrust::iterator_traits::difference_type + unique_count(my_tag, + ForwardIterator, + ForwardIterator) +{ + return 13; +} + +void TestUniqueCountDispatchImplicit() +{ + thrust::device_vector vec(1); + + auto result = thrust::unique_count( + thrust::retag(vec.begin()), + thrust::retag(vec.begin())); + + ASSERT_EQUAL(13, result); +} +DECLARE_UNITTEST(TestUniqueCountDispatchImplicit); + + template struct is_equal_div_10_unique { @@ -266,3 +310,48 @@ struct TestUniqueCopyToDiscardIterator VariableUnitTest TestUniqueCopyToDiscardIteratorInstance; +template +void TestUniqueCountSimple(void) +{ + typedef typename Vector::value_type T; + + 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; + + int count = thrust::unique_count(data.begin(), data.end()); + + ASSERT_EQUAL(count, 7); + + int div_10_count = thrust::unique_count(data.begin(), data.end(), is_equal_div_10_unique()); + + ASSERT_EQUAL(div_10_count, 3); +} +DECLARE_INTEGRAL_VECTOR_UNITTEST(TestUniqueCountSimple); + +template +struct TestUniqueCount +{ + void operator()(const size_t n) + { + thrust::host_vector h_data = unittest::random_integers(n); + thrust::device_vector d_data = h_data; + + int h_count{}; + int d_count{}; + + h_count = thrust::unique_count(h_data.begin(), h_data.end()); + d_count = thrust::unique_count(d_data.begin(), d_data.end()); + + ASSERT_EQUAL(h_count, d_count); + } +}; +VariableUnitTest TestUniqueCountInstance; diff --git a/thrust/count.h b/thrust/count.h index 52b22d205..abf8b2d6c 100644 --- a/thrust/count.h +++ b/thrust/count.h @@ -228,4 +228,4 @@ template THRUST_NAMESPACE_END -#include +#include diff --git a/thrust/detail/count.h b/thrust/detail/count.h new file mode 100644 index 000000000..7c48bc546 --- /dev/null +++ b/thrust/detail/count.h @@ -0,0 +1,60 @@ +/* + * 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 +#include + +THRUST_NAMESPACE_BEGIN + +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + count(const thrust::detail::execution_policy_base &exec, + InputIterator first, + InputIterator last, + const EqualityComparable& value); + +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + count_if(const thrust::detail::execution_policy_base &exec, + InputIterator first, + InputIterator last, + Predicate pred); + +template + typename thrust::iterator_traits::difference_type + count(InputIterator first, + InputIterator last, + const EqualityComparable& value); + +template + typename thrust::iterator_traits::difference_type + count_if(InputIterator first, + InputIterator last, + Predicate pred); + +THRUST_NAMESPACE_END + +#include diff --git a/thrust/detail/unique.inl b/thrust/detail/unique.inl index a1a7b492b..ac5475f02 100644 --- a/thrust/detail/unique.inl +++ b/thrust/detail/unique.inl @@ -327,6 +327,67 @@ template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(const thrust::detail::execution_policy_base &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred) +{ + using thrust::system::detail::generic::unique_count; + return unique_count(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last, binary_pred); +} // end unique_count() + +__thrust_exec_check_disable__ +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(const thrust::detail::execution_policy_base &exec, + ForwardIterator first, + ForwardIterator last) +{ + using thrust::system::detail::generic::unique_count; + return unique_count(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first, last); +} // end unique_count() + +__thrust_exec_check_disable__ +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred) +{ + using thrust::system::detail::generic::select_system; + + typedef typename thrust::iterator_system::type System; + + System system; + + return thrust::unique_count(select_system(system), first, last, binary_pred); +} // end unique_count() + +__thrust_exec_check_disable__ +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(ForwardIterator first, + ForwardIterator last) +{ + using thrust::system::detail::generic::select_system; + + typedef typename thrust::iterator_system::type System; + + System system; + + return thrust::unique_count(select_system(system), first, last); +} // end unique_count() THRUST_NAMESPACE_END diff --git a/thrust/system/cuda/detail/unique.h b/thrust/system/cuda/detail/unique.h index 91dd2b84f..603c5e3b2 100644 --- a/thrust/system/cuda/detail/unique.h +++ b/thrust/system/cuda/detail/unique.h @@ -69,6 +69,16 @@ unique_copy( OutputIterator result, BinaryPredicate binary_pred); +template +__host__ __device__ typename thrust::iterator_traits::difference_type +unique_count( + const thrust::detail::execution_policy_base &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred); + namespace cuda_cub { // XXX it should be possible to unify unique & unique_by_key into a single @@ -794,6 +804,36 @@ unique(execution_policy &policy, return cuda_cub::unique(policy, first, last, equal_to()); } + +template +struct zip_adj_not_predicate { + template + bool __host__ __device__ operator()(TupleType&& tuple) { + return !binary_pred(thrust::get<0>(tuple), thrust::get<1>(tuple)); + } + + BinaryPred binary_pred; +}; + + +__thrust_exec_check_disable__ +template +typename thrust::iterator_traits::difference_type __host__ __device__ +unique_count(execution_policy &policy, + InputIt first, + InputIt last, + BinaryPred binary_pred) +{ + if (first == last) { + return 0; + } + auto size = last - first; + auto it = thrust::make_zip_iterator(thrust::make_tuple(first, first + 1)); + return 1 + thrust::count_if(policy, it, it + (size - 1), zip_adj_not_predicate{binary_pred}); +} + } // namespace cuda_cub THRUST_NAMESPACE_END diff --git a/thrust/system/detail/generic/unique.h b/thrust/system/detail/generic/unique.h index 5f008978f..ce3bff884 100644 --- a/thrust/system/detail/generic/unique.h +++ b/thrust/system/detail/generic/unique.h @@ -68,6 +68,26 @@ OutputIterator unique_copy(thrust::execution_policy &exec, BinaryPredicate binary_pred); +template +__host__ __device__ +typename thrust::iterator_traits::difference_type + unique_count(thrust::execution_policy &exec, + ForwardIterator first, + ForwardIterator last); + + +template +__host__ __device__ +typename thrust::iterator_traits::difference_type + unique_count(thrust::execution_policy &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred); + + } // end namespace generic } // end namespace detail } // end namespace system diff --git a/thrust/system/detail/generic/unique.inl b/thrust/system/detail/generic/unique.inl index 5d3ba2fd1..bb66e3585 100644 --- a/thrust/system/detail/generic/unique.inl +++ b/thrust/system/detail/generic/unique.inl @@ -24,6 +24,7 @@ #include #include #include +#include #include #include #include @@ -100,6 +101,37 @@ __host__ __device__ } // end unique_copy() +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(thrust::execution_policy &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred) +{ + thrust::detail::head_flags stencil(first, last, binary_pred); + + using namespace thrust::placeholders; + + return thrust::count_if(exec, stencil.begin(), stencil.end(), _1); +} // end unique_copy() + + +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(thrust::execution_policy &exec, + ForwardIterator first, + ForwardIterator last) +{ + typedef typename thrust::iterator_value::type value_type; + return thrust::unique_count(exec, first, last, thrust::equal_to()); +} // end unique_copy() + + } // end namespace generic } // end namespace detail } // end namespace system diff --git a/thrust/system/detail/sequential/unique.h b/thrust/system/detail/sequential/unique.h index e4953e9ae..c4fe5268a 100644 --- a/thrust/system/detail/sequential/unique.h +++ b/thrust/system/detail/sequential/unique.h @@ -89,6 +89,40 @@ __host__ __device__ } // end unique() +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(sequential::execution_policy &, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred) +{ + typedef typename thrust::iterator_traits::value_type T; + typename thrust::iterator_traits::difference_type count{}; + + if(first != last) + { + count++; + T prev = *first; + + for(++first; first != last; ++first) + { + T temp = *first; + + if (!binary_pred(prev, temp)) + { + count++; + prev = temp; + } + } + } + + return count; +} // end unique() + + } // end namespace sequential } // end namespace detail } // end namespace system diff --git a/thrust/system/omp/detail/unique.h b/thrust/system/omp/detail/unique.h index 304caf66d..cf8025665 100644 --- a/thrust/system/omp/detail/unique.h +++ b/thrust/system/omp/detail/unique.h @@ -49,6 +49,16 @@ template + typename thrust::iterator_traits::difference_type + unique_count(execution_policy &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred); + + } // end namespace detail } // end namespace omp } // end namespace system diff --git a/thrust/system/omp/detail/unique.inl b/thrust/system/omp/detail/unique.inl index c03203efe..5425668e7 100644 --- a/thrust/system/omp/detail/unique.inl +++ b/thrust/system/omp/detail/unique.inl @@ -58,6 +58,21 @@ template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(execution_policy &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred) +{ + // omp prefers generic::unique_count to cpp::unique_count + return thrust::system::detail::generic::unique_count(exec,first,last,binary_pred); +} // end unique_count() + + } // end namespace detail } // end namespace omp } // end namespace system diff --git a/thrust/system/tbb/detail/unique.h b/thrust/system/tbb/detail/unique.h index db4692d34..843e6406e 100644 --- a/thrust/system/tbb/detail/unique.h +++ b/thrust/system/tbb/detail/unique.h @@ -49,6 +49,16 @@ template + typename thrust::iterator_traits::difference_type + unique_count(execution_policy &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred); + + } // end namespace detail } // end namespace tbb } // end namespace system diff --git a/thrust/system/tbb/detail/unique.inl b/thrust/system/tbb/detail/unique.inl index 0c3c16f2e..4a3b0b332 100644 --- a/thrust/system/tbb/detail/unique.inl +++ b/thrust/system/tbb/detail/unique.inl @@ -58,6 +58,21 @@ template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(execution_policy &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred) +{ + // omp prefers generic::unique_count to cpp::unique_count + return thrust::system::detail::generic::unique_count(exec,first,last,binary_pred); +} // end unique_count() + + } // end namespace detail } // end namespace tbb } // end namespace system diff --git a/thrust/unique.h b/thrust/unique.h index 426b37ab7..234cd4935 100644 --- a/thrust/unique.h +++ b/thrust/unique.h @@ -23,6 +23,7 @@ #include #include +#include #include THRUST_NAMESPACE_BEGIN @@ -956,6 +957,180 @@ template[first, last) + * with the same value, + * + * This version of \p unique_count uses the function object \p binary_pred to test for equality. + * + * The algorithm's execution is parallelized as determined by \p exec. + * + * \param exec The execution policy to use for parallelization. + * \param first The beginning of the input range. + * \param last The end of the input range. + * \param binary_pred The binary predicate used to determine equality. + * \return The number of runs of equal elements in [first, new_last) + * + * \tparam DerivedPolicy The name of the derived execution policy. + * \tparam ForwardIterator is a model of Forward Iterator, + * and \p ForwardIterator's \c value_type is convertible to \p BinaryPredicate's \c first_argument_type and to \p BinaryPredicate's \c second_argument_type. + * \tparam BinaryPredicate is a model of Binary Predicate. + * + * The following code snippet demonstrates how to use \p unique_count to + * determine a number of runs of equal elements using the \p thrust::host execution policy + * for parallelization: + * + * \code + * #include + * #include + * ... + * const int N = 7; + * int A[N] = {1, 3, 3, 3, 2, 2, 1}; + * int count = thrust::unique_count(thrust::host, A, A + N, thrust::equal_to()); + * // count is now 4 + * \endcode + * + * \see unique_copy + * \see unique_by_key_copy + * \see reduce_by_key_copy + */ +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(const thrust::detail::execution_policy_base &exec, + ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred); + + +/*! \p unique_count counts runs of equal elements in the range [first, last) + * with the same value, + * + * This version of \p unique_count uses \c operator== to test for equality. + * + * The algorithm's execution is parallelized as determined by \p exec. + * + * \param exec The execution policy to use for parallelization. + * \param first The beginning of the input range. + * \param last The end of the input range. + * \param binary_pred The binary predicate used to determine equality. + * \return The number of runs of equal elements in [first, new_last) + * + * \tparam DerivedPolicy The name of the derived execution policy. + * \tparam ForwardIterator is a model of Forward Iterator, + * and \p ForwardIterator's \c value_type is convertible to \p BinaryPredicate's \c first_argument_type and to \p BinaryPredicate's \c second_argument_type. + * \tparam BinaryPredicate is a model of Binary Predicate. + * + * The following code snippet demonstrates how to use \p unique_count to + * determine the number of runs of equal elements using the \p thrust::host execution policy + * for parallelization: + * + * \code + * #include + * #include + * ... + * const int N = 7; + * int A[N] = {1, 3, 3, 3, 2, 2, 1}; + * int count = thrust::unique_count(thrust::host, A, A + N); + * // count is now 4 + * \endcode + * + * \see unique_copy + * \see unique_by_key_copy + * \see reduce_by_key_copy + */ +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(const thrust::detail::execution_policy_base &exec, + ForwardIterator first, + ForwardIterator last); + + +/*! \p unique_count counts runs of equal elements in the range [first, last) + * with the same value, + * + * This version of \p unique_count uses the function object \p binary_pred to test for equality. + * + * \param exec The execution policy to use for parallelization. + * \param first The beginning of the input range. + * \param last The end of the input range. + * \param binary_pred The binary predicate used to determine equality. + * \return The number of runs of equal elements in [first, new_last) + * + * \tparam DerivedPolicy The name of the derived execution policy. + * \tparam ForwardIterator is a model of Forward Iterator, + * and \p ForwardIterator's \c value_type is convertible to \p BinaryPredicate's \c first_argument_type and to \p BinaryPredicate's \c second_argument_type. + * \tparam BinaryPredicate is a model of Binary Predicate. + * + * The following code snippet demonstrates how to use \p unique_count to + * determine the number of runs of equal elements: + * + * \code + * #include + * #include + * ... + * const int N = 7; + * int A[N] = {1, 3, 3, 3, 2, 2, 1}; + * int count = thrust::unique_count(A, A + N, thrust::equal_to()); + * // count is now 4 + * \endcode + * + * \see unique_copy + * \see unique_by_key_copy + * \see reduce_by_key_copy + */ +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(ForwardIterator first, + ForwardIterator last, + BinaryPredicate binary_pred); + + +/*! \p unique_count counts runs of equal elements in the range [first, last) + * with the same value, + * + * This version of \p unique_count uses \c operator== to test for equality. + * + * \param exec The execution policy to use for parallelization. + * \param first The beginning of the input range. + * \param last The end of the input range. + * \param binary_pred The binary predicate used to determine equality. + * \return The number of runs of equal elements in [first, new_last) + * + * \tparam DerivedPolicy The name of the derived execution policy. + * \tparam ForwardIterator is a model of Forward Iterator, + * and \p ForwardIterator's \c value_type is convertible to \p BinaryPredicate's \c first_argument_type and to \p BinaryPredicate's \c second_argument_type. + * \tparam BinaryPredicate is a model of Binary Predicate. + * + * The following code snippet demonstrates how to use \p unique_count to + * determine the number of runs of equal elements: + * + * \code + * #include + * #include + * ... + * const int N = 7; + * int A[N] = {1, 3, 3, 3, 2, 2, 1}; + * int count = thrust::unique_count(thrust::host, A, A + N); + * // count is now 4 + * \endcode + * + * \see unique_copy + * \see unique_by_key_copy + * \see reduce_by_key_copy + */ +template +__host__ __device__ + typename thrust::iterator_traits::difference_type + unique_count(ForwardIterator first, + ForwardIterator last); + + /*! \} // end stream_compaction */ From 0c354e89d63d5b3803d0c09e4b3fae4c7d9e83e5 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Wed, 2 Mar 2022 12:55:29 +0100 Subject: [PATCH 02/10] unique_count: weaken iterator requirements --- thrust/system/cuda/detail/unique.h | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/thrust/system/cuda/detail/unique.h b/thrust/system/cuda/detail/unique.h index 603c5e3b2..83d9f058c 100644 --- a/thrust/system/cuda/detail/unique.h +++ b/thrust/system/cuda/detail/unique.h @@ -41,6 +41,7 @@ #include #include #include +#include #include #include @@ -818,20 +819,20 @@ struct zip_adj_not_predicate { __thrust_exec_check_disable__ template -typename thrust::iterator_traits::difference_type __host__ __device__ +typename thrust::iterator_traits::difference_type __host__ __device__ unique_count(execution_policy &policy, - InputIt first, - InputIt last, + ForwardIt first, + ForwardIt last, BinaryPred binary_pred) { if (first == last) { return 0; } - auto size = last - first; - auto it = thrust::make_zip_iterator(thrust::make_tuple(first, first + 1)); - return 1 + thrust::count_if(policy, it, it + (size - 1), zip_adj_not_predicate{binary_pred}); + auto size = thrust::distance(first, last); + auto it = thrust::make_zip_iterator(thrust::make_tuple(first, thrust::next(first))); + return 1 + thrust::count_if(policy, it, thrust::next(it, size - 1), zip_adj_not_predicate{binary_pred}); } } // namespace cuda_cub From fca96ec32eaa72d9f944c10101198cc45499a5b3 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Fri, 11 Mar 2022 15:52:16 +0100 Subject: [PATCH 03/10] unique: improve template parameter naming The interface specifies ForwardIterator, not InputIterator --- thrust/system/cuda/detail/unique.h | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/thrust/system/cuda/detail/unique.h b/thrust/system/cuda/detail/unique.h index 83d9f058c..89f1ea76e 100644 --- a/thrust/system/cuda/detail/unique.h +++ b/thrust/system/cuda/detail/unique.h @@ -769,15 +769,15 @@ unique_copy(execution_policy &policy, __thrust_exec_check_disable__ template -InputIt __host__ __device__ +ForwardIt __host__ __device__ unique(execution_policy &policy, - InputIt first, - InputIt last, + ForwardIt first, + ForwardIt last, BinaryPred binary_pred) { - InputIt ret = first; + ForwardIt ret = first; if (__THRUST_HAS_CUDART__) { ret = cuda_cub::unique_copy(policy, first, last, first, binary_pred); @@ -795,13 +795,13 @@ unique(execution_policy &policy, } template -InputIt __host__ __device__ + class ForwardIt> +ForwardIt __host__ __device__ unique(execution_policy &policy, - InputIt first, - InputIt last) + ForwardIt first, + ForwardIt last) { - typedef typename iterator_traits::value_type input_type; + typedef typename iterator_traits::value_type input_type; return cuda_cub::unique(policy, first, last, equal_to()); } From 0b41e08165825d55145442ebe07e87c3dc85351f Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Fri, 11 Mar 2022 17:08:04 +0100 Subject: [PATCH 04/10] unique: test with ForwardIterator parameters --- testing/unique.cu | 22 ++++++++---- testing/unittest/iterator_helpers.h | 52 +++++++++++++++++++++++++++++ 2 files changed, 67 insertions(+), 7 deletions(-) create mode 100644 testing/unittest/iterator_helpers.h diff --git a/testing/unique.cu b/testing/unique.cu index 7df2def87..b0ae8dec0 100644 --- a/testing/unique.cu +++ b/testing/unique.cu @@ -1,4 +1,5 @@ #include +#include #include #include #include @@ -163,11 +164,13 @@ void TestUniqueSimple(void) data[8] = 31; data[9] = 37; - typename Vector::iterator new_last; + forward_iterator_wrapper new_last; + const auto begin = make_forward_iterator_wrapper(data.begin()); + const auto end = make_forward_iterator_wrapper(data.end()); - new_last = thrust::unique(data.begin(), data.end()); + new_last = thrust::unique(begin, end); - ASSERT_EQUAL(new_last - data.begin(), 7); + ASSERT_EQUAL(thrust::distance(begin, new_last), 7); ASSERT_EQUAL(data[0], 11); ASSERT_EQUAL(data[1], 12); ASSERT_EQUAL(data[2], 20); @@ -176,9 +179,9 @@ void TestUniqueSimple(void) ASSERT_EQUAL(data[5], 31); ASSERT_EQUAL(data[6], 37); - new_last = thrust::unique(data.begin(), new_last, is_equal_div_10_unique()); + new_last = thrust::unique(begin, new_last, is_equal_div_10_unique()); - ASSERT_EQUAL(new_last - data.begin(), 3); + ASSERT_EQUAL(thrust::distance(begin, new_last), 3); ASSERT_EQUAL(data[0], 11); ASSERT_EQUAL(data[1], 20); ASSERT_EQUAL(data[2], 31); @@ -327,11 +330,16 @@ void TestUniqueCountSimple(void) data[8] = 31; data[9] = 37; - int count = thrust::unique_count(data.begin(), data.end()); + int count = thrust::unique_count( + make_forward_iterator_wrapper(data.begin()), + make_forward_iterator_wrapper(data.end())); ASSERT_EQUAL(count, 7); - int div_10_count = thrust::unique_count(data.begin(), data.end(), is_equal_div_10_unique()); + int div_10_count = thrust::unique_count( + make_forward_iterator_wrapper(data.begin()), + make_forward_iterator_wrapper(data.end()), + is_equal_div_10_unique()); ASSERT_EQUAL(div_10_count, 3); } diff --git a/testing/unittest/iterator_helpers.h b/testing/unittest/iterator_helpers.h new file mode 100644 index 000000000..f6ac00339 --- /dev/null +++ b/testing/unittest/iterator_helpers.h @@ -0,0 +1,52 @@ +#pragma once + +#include + + +// Wraps an existing iterator into a forward iterator, +// thus removing some of its functionality +template +struct forward_iterator_wrapper { + // LegacyIterator requirements + using reference = typename Iterator::reference; + using pointer = typename Iterator::pointer; + using value_type = typename Iterator::value_type; + using difference_type = typename Iterator::difference_type; + using iterator_category = std::forward_iterator_tag; + + __host__ __device__ reference operator*() const { + return *wrapped; + } + + __host__ __device__ forward_iterator_wrapper& operator++() { + ++wrapped; + return *this; + } + + // LegacyInputIterator + __host__ __device__ bool operator==(const forward_iterator_wrapper& other) { + return wrapped == other.wrapped; + } + + __host__ __device__ bool operator!=(const forward_iterator_wrapper& other) { + return !(*this == other); + } + + __host__ __device__ forward_iterator_wrapper operator++(int) { + auto cpy = *this; + ++(*this); + return cpy; + } + + __host__ __device__ pointer operator->() const { + return wrapped.operator->(); + } + + Iterator wrapped; +}; + + +template +forward_iterator_wrapper make_forward_iterator_wrapper(Iterator it) { + return {it}; +} From 1532df8007ff38189cdb88738eafb1759b90b377 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Fri, 25 Mar 2022 22:37:18 +0100 Subject: [PATCH 05/10] improve forward_iterator_wrapper * use iterator traits * use hidden friend operators * fix member access operator Co-authored-by: Eric Niebler --- testing/unittest/iterator_helpers.h | 32 +++++++++++++++++++---------- 1 file changed, 21 insertions(+), 11 deletions(-) diff --git a/testing/unittest/iterator_helpers.h b/testing/unittest/iterator_helpers.h index f6ac00339..8d4f03f56 100644 --- a/testing/unittest/iterator_helpers.h +++ b/testing/unittest/iterator_helpers.h @@ -1,6 +1,6 @@ #pragma once -#include +#include // Wraps an existing iterator into a forward iterator, @@ -8,11 +8,15 @@ template struct forward_iterator_wrapper { // LegacyIterator requirements - using reference = typename Iterator::reference; - using pointer = typename Iterator::pointer; - using value_type = typename Iterator::value_type; - using difference_type = typename Iterator::difference_type; + using reference = typename thrust::iterator_traits::reference; + using pointer = typename thrust::iterator_traits::pointer; + using value_type = typename thrust::iterator_traits::value_type; + using difference_type = typename thrust::iterator_traits::difference_type; using iterator_category = std::forward_iterator_tag; + using base_iterator_category = typename thrust::iterator_traits::iterator_category; + static_assert( + std::is_convertible::value, + "Cannot create forward_iterator_wrapper around an iterator that is not itself at least a forward iterator"); __host__ __device__ reference operator*() const { return *wrapped; @@ -24,12 +28,12 @@ struct forward_iterator_wrapper { } // LegacyInputIterator - __host__ __device__ bool operator==(const forward_iterator_wrapper& other) { - return wrapped == other.wrapped; + friend __host__ __device__ bool operator==(const forward_iterator_wrapper& a, const forward_iterator_wrapper& b) { + return a.wrapped == b.wrapped; } - __host__ __device__ bool operator!=(const forward_iterator_wrapper& other) { - return !(*this == other); + friend __host__ __device__ bool operator!=(const forward_iterator_wrapper& a, const forward_iterator_wrapper& b) { + return !(a == b); } __host__ __device__ forward_iterator_wrapper operator++(int) { @@ -37,8 +41,14 @@ struct forward_iterator_wrapper { ++(*this); return cpy; } - - __host__ __device__ pointer operator->() const { + + template + __host__ __device__ typename std::enable_if::value, pointer>::type operator->() const { + return wrapped; + } + + template + __host__ __device__ typename std::enable_if::value, pointer>::type operator->() const { return wrapped.operator->(); } From e37433cb65564e19fcd2acce20cdd4225b7fa256 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Fri, 25 Mar 2022 22:37:52 +0100 Subject: [PATCH 06/10] unique_count: add missing cuda tests --- testing/cuda/unique.cu | 125 +++++++++++++++++++++++++++++++++++++++++ 1 file changed, 125 insertions(+) diff --git a/testing/cuda/unique.cu b/testing/cuda/unique.cu index 3e404238f..2fef6b61f 100644 --- a/testing/cuda/unique.cu +++ b/testing/cuda/unique.cu @@ -320,3 +320,128 @@ void TestUniqueCopyCudaStreamsNoSync() } DECLARE_UNITTEST(TestUniqueCopyCudaStreamsNoSync); + +template +__global__ +void unique_count_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 result) +{ + *result = thrust::unique_count(exec, first, last); +} + + +template +__global__ +void unique_count_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, BinaryPredicate pred, Iterator2 result) +{ + *result = thrust::unique_count(exec, first, last, pred); +} + + +template +void TestUniqueCountDevice(ExecutionPolicy exec) +{ + typedef thrust::device_vector Vector; + typedef Vector::value_type T; + + 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(1, -1); + + unique_count_kernel<<<1,1>>>(exec, data.begin(), data.end(), output.begin()); + { + cudaError_t const err = cudaDeviceSynchronize(); + ASSERT_EQUAL(cudaSuccess, err); + } + + ASSERT_EQUAL(output[0], 7); + + unique_count_kernel<<<1,1>>>(exec, data.begin(), data.end(), is_equal_div_10_unique(), output.begin()); + { + cudaError_t const err = cudaDeviceSynchronize(); + ASSERT_EQUAL(cudaSuccess, err); + } + + ASSERT_EQUAL(output[0], 3); +} + + +void TestUniqueCountDeviceSeq() +{ + TestUniqueCountDevice(thrust::seq); +} +DECLARE_UNITTEST(TestUniqueCountDeviceSeq); + + +void TestUniqueCountDeviceDevice() +{ + TestUniqueCountDevice(thrust::device); +} +DECLARE_UNITTEST(TestUniqueCountDeviceDevice); + + +void TestUniqueCountDeviceNoSync() +{ + TestUniqueCountDevice(thrust::cuda::par_nosync); +} +DECLARE_UNITTEST(TestUniqueCountDeviceNoSync); + + +template +void TestUniqueCountCudaStreams(ExecutionPolicy policy) +{ + typedef thrust::device_vector Vector; + typedef Vector::value_type T; + + 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; + + cudaStream_t s; + cudaStreamCreate(&s); + + auto streampolicy = policy.on(s); + + int result = thrust::unique_count(streampolicy, data.begin(), data.end()); + cudaStreamSynchronize(s); + + ASSERT_EQUAL(result, 7); + + result = thrust::unique_count(streampolicy, data.begin(), data.end(), is_equal_div_10_unique()); + cudaStreamSynchronize(s); + + ASSERT_EQUAL(result, 3); + + cudaStreamDestroy(s); +} + +void TestUniqueCountCudaStreamsSync() +{ + TestUniqueCountCudaStreams(thrust::cuda::par); +} +DECLARE_UNITTEST(TestUniqueCountCudaStreamsSync); + + +void TestUniqueCountCudaStreamsNoSync() +{ + TestUniqueCountCudaStreams(thrust::cuda::par_nosync); +} +DECLARE_UNITTEST(TestUniqueCountCudaStreamsNoSync); + From fac36573bec54519d713d06f46fa45292714a7d2 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Thu, 5 May 2022 22:12:49 +0200 Subject: [PATCH 07/10] use thrust iterator categories in iterator wrapper --- testing/unittest/iterator_helpers.h | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/testing/unittest/iterator_helpers.h b/testing/unittest/iterator_helpers.h index 8d4f03f56..4b5809b38 100644 --- a/testing/unittest/iterator_helpers.h +++ b/testing/unittest/iterator_helpers.h @@ -1,6 +1,8 @@ #pragma once #include +#include +#include // Wraps an existing iterator into a forward iterator, @@ -8,14 +10,21 @@ template struct forward_iterator_wrapper { // LegacyIterator requirements + using iterator_system_tag = typename thrust::iterator_system::type; using reference = typename thrust::iterator_traits::reference; using pointer = typename thrust::iterator_traits::pointer; using value_type = typename thrust::iterator_traits::value_type; using difference_type = typename thrust::iterator_traits::difference_type; - using iterator_category = std::forward_iterator_tag; + using iterator_category = typename std::conditional< + std::is_convertible::value, + thrust::forward_device_iterator_tag, + typename std::conditional< + std::is_convertible::value, + thrust::forward_host_iterator_tag, + std::forward_iterator_tag>::type>::type; using base_iterator_category = typename thrust::iterator_traits::iterator_category; static_assert( - std::is_convertible::value, + std::is_convertible::value, "Cannot create forward_iterator_wrapper around an iterator that is not itself at least a forward iterator"); __host__ __device__ reference operator*() const { From a865d5350bdaab7efeb9cd4c56b023485b9d77e4 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Fri, 6 May 2022 06:40:28 +0200 Subject: [PATCH 08/10] Revert "use thrust iterator categories in iterator wrapper" This reverts commit fac36573bec54519d713d06f46fa45292714a7d2. --- testing/unittest/iterator_helpers.h | 13 ++----------- 1 file changed, 2 insertions(+), 11 deletions(-) diff --git a/testing/unittest/iterator_helpers.h b/testing/unittest/iterator_helpers.h index 4b5809b38..8d4f03f56 100644 --- a/testing/unittest/iterator_helpers.h +++ b/testing/unittest/iterator_helpers.h @@ -1,8 +1,6 @@ #pragma once #include -#include -#include // Wraps an existing iterator into a forward iterator, @@ -10,21 +8,14 @@ template struct forward_iterator_wrapper { // LegacyIterator requirements - using iterator_system_tag = typename thrust::iterator_system::type; using reference = typename thrust::iterator_traits::reference; using pointer = typename thrust::iterator_traits::pointer; using value_type = typename thrust::iterator_traits::value_type; using difference_type = typename thrust::iterator_traits::difference_type; - using iterator_category = typename std::conditional< - std::is_convertible::value, - thrust::forward_device_iterator_tag, - typename std::conditional< - std::is_convertible::value, - thrust::forward_host_iterator_tag, - std::forward_iterator_tag>::type>::type; + using iterator_category = std::forward_iterator_tag; using base_iterator_category = typename thrust::iterator_traits::iterator_category; static_assert( - std::is_convertible::value, + std::is_convertible::value, "Cannot create forward_iterator_wrapper around an iterator that is not itself at least a forward iterator"); __host__ __device__ reference operator*() const { From 7bf9735a0bdf8199ff35dbeb565bf5b1f5290bdd Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Fri, 6 May 2022 06:40:32 +0200 Subject: [PATCH 09/10] Revert "improve forward_iterator_wrapper" This reverts commit 1532df8007ff38189cdb88738eafb1759b90b377. --- testing/unittest/iterator_helpers.h | 32 ++++++++++------------------- 1 file changed, 11 insertions(+), 21 deletions(-) diff --git a/testing/unittest/iterator_helpers.h b/testing/unittest/iterator_helpers.h index 8d4f03f56..f6ac00339 100644 --- a/testing/unittest/iterator_helpers.h +++ b/testing/unittest/iterator_helpers.h @@ -1,6 +1,6 @@ #pragma once -#include +#include // Wraps an existing iterator into a forward iterator, @@ -8,15 +8,11 @@ template struct forward_iterator_wrapper { // LegacyIterator requirements - using reference = typename thrust::iterator_traits::reference; - using pointer = typename thrust::iterator_traits::pointer; - using value_type = typename thrust::iterator_traits::value_type; - using difference_type = typename thrust::iterator_traits::difference_type; + using reference = typename Iterator::reference; + using pointer = typename Iterator::pointer; + using value_type = typename Iterator::value_type; + using difference_type = typename Iterator::difference_type; using iterator_category = std::forward_iterator_tag; - using base_iterator_category = typename thrust::iterator_traits::iterator_category; - static_assert( - std::is_convertible::value, - "Cannot create forward_iterator_wrapper around an iterator that is not itself at least a forward iterator"); __host__ __device__ reference operator*() const { return *wrapped; @@ -28,12 +24,12 @@ struct forward_iterator_wrapper { } // LegacyInputIterator - friend __host__ __device__ bool operator==(const forward_iterator_wrapper& a, const forward_iterator_wrapper& b) { - return a.wrapped == b.wrapped; + __host__ __device__ bool operator==(const forward_iterator_wrapper& other) { + return wrapped == other.wrapped; } - friend __host__ __device__ bool operator!=(const forward_iterator_wrapper& a, const forward_iterator_wrapper& b) { - return !(a == b); + __host__ __device__ bool operator!=(const forward_iterator_wrapper& other) { + return !(*this == other); } __host__ __device__ forward_iterator_wrapper operator++(int) { @@ -41,14 +37,8 @@ struct forward_iterator_wrapper { ++(*this); return cpy; } - - template - __host__ __device__ typename std::enable_if::value, pointer>::type operator->() const { - return wrapped; - } - - template - __host__ __device__ typename std::enable_if::value, pointer>::type operator->() const { + + __host__ __device__ pointer operator->() const { return wrapped.operator->(); } From 57f8e5e15beb27b23b0e133d73ec7b499d46157d Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Fri, 6 May 2022 06:40:36 +0200 Subject: [PATCH 10/10] Revert "unique: test with ForwardIterator parameters" This reverts commit 0b41e08165825d55145442ebe07e87c3dc85351f. --- testing/unique.cu | 22 ++++-------- testing/unittest/iterator_helpers.h | 52 ----------------------------- 2 files changed, 7 insertions(+), 67 deletions(-) delete mode 100644 testing/unittest/iterator_helpers.h diff --git a/testing/unique.cu b/testing/unique.cu index b0ae8dec0..7df2def87 100644 --- a/testing/unique.cu +++ b/testing/unique.cu @@ -1,5 +1,4 @@ #include -#include #include #include #include @@ -164,13 +163,11 @@ void TestUniqueSimple(void) data[8] = 31; data[9] = 37; - forward_iterator_wrapper new_last; - const auto begin = make_forward_iterator_wrapper(data.begin()); - const auto end = make_forward_iterator_wrapper(data.end()); + typename Vector::iterator new_last; - new_last = thrust::unique(begin, end); + new_last = thrust::unique(data.begin(), data.end()); - ASSERT_EQUAL(thrust::distance(begin, new_last), 7); + ASSERT_EQUAL(new_last - data.begin(), 7); ASSERT_EQUAL(data[0], 11); ASSERT_EQUAL(data[1], 12); ASSERT_EQUAL(data[2], 20); @@ -179,9 +176,9 @@ void TestUniqueSimple(void) ASSERT_EQUAL(data[5], 31); ASSERT_EQUAL(data[6], 37); - new_last = thrust::unique(begin, new_last, is_equal_div_10_unique()); + new_last = thrust::unique(data.begin(), new_last, is_equal_div_10_unique()); - ASSERT_EQUAL(thrust::distance(begin, new_last), 3); + ASSERT_EQUAL(new_last - data.begin(), 3); ASSERT_EQUAL(data[0], 11); ASSERT_EQUAL(data[1], 20); ASSERT_EQUAL(data[2], 31); @@ -330,16 +327,11 @@ void TestUniqueCountSimple(void) data[8] = 31; data[9] = 37; - int count = thrust::unique_count( - make_forward_iterator_wrapper(data.begin()), - make_forward_iterator_wrapper(data.end())); + int count = thrust::unique_count(data.begin(), data.end()); ASSERT_EQUAL(count, 7); - int div_10_count = thrust::unique_count( - make_forward_iterator_wrapper(data.begin()), - make_forward_iterator_wrapper(data.end()), - is_equal_div_10_unique()); + int div_10_count = thrust::unique_count(data.begin(), data.end(), is_equal_div_10_unique()); ASSERT_EQUAL(div_10_count, 3); } diff --git a/testing/unittest/iterator_helpers.h b/testing/unittest/iterator_helpers.h deleted file mode 100644 index f6ac00339..000000000 --- a/testing/unittest/iterator_helpers.h +++ /dev/null @@ -1,52 +0,0 @@ -#pragma once - -#include - - -// Wraps an existing iterator into a forward iterator, -// thus removing some of its functionality -template -struct forward_iterator_wrapper { - // LegacyIterator requirements - using reference = typename Iterator::reference; - using pointer = typename Iterator::pointer; - using value_type = typename Iterator::value_type; - using difference_type = typename Iterator::difference_type; - using iterator_category = std::forward_iterator_tag; - - __host__ __device__ reference operator*() const { - return *wrapped; - } - - __host__ __device__ forward_iterator_wrapper& operator++() { - ++wrapped; - return *this; - } - - // LegacyInputIterator - __host__ __device__ bool operator==(const forward_iterator_wrapper& other) { - return wrapped == other.wrapped; - } - - __host__ __device__ bool operator!=(const forward_iterator_wrapper& other) { - return !(*this == other); - } - - __host__ __device__ forward_iterator_wrapper operator++(int) { - auto cpy = *this; - ++(*this); - return cpy; - } - - __host__ __device__ pointer operator->() const { - return wrapped.operator->(); - } - - Iterator wrapped; -}; - - -template -forward_iterator_wrapper make_forward_iterator_wrapper(Iterator it) { - return {it}; -}