From b9dea3d97362da8f68ed49309e5d2518987b8164 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 19 Feb 2025 16:18:49 +0100 Subject: [PATCH 1/4] Add tests for counting_iterator difference_type --- thrust/testing/counting_iterator.cu | 28 ++++++++++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/thrust/testing/counting_iterator.cu b/thrust/testing/counting_iterator.cu index a7fca340523..90abe1e4bc8 100644 --- a/thrust/testing/counting_iterator.cu +++ b/thrust/testing/counting_iterator.cu @@ -10,6 +10,34 @@ #include +template +inline constexpr bool diff_type_is = + ::cuda::std::is_same_v::difference_type, DifferenceType>; + +static_assert(diff_type_is); +static_assert(diff_type_is); +static_assert(diff_type_is); +static_assert(diff_type_is); +static_assert(diff_type_is); +static_assert(diff_type_is); +static_assert(diff_type_is); +static_assert(diff_type_is); +#if _CCCL_HAS_INT128() +static_assert(diff_type_is<__int128_t, __int128_t>); +static_assert(diff_type_is<__uint128_t, long>); +#endif +static_assert(diff_type_is); +static_assert(diff_type_is); + +struct custom_int +{ + _CCCL_HOST_DEVICE custom_int(int) {} + _CCCL_HOST_DEVICE operator int() const; +}; +static_assert(thrust::detail::is_numeric::value); + +static_assert(diff_type_is); + _CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_MSVC(4244 4267) // possible loss of data From 22d84f483171c78f60895f31f089d66b03179a69 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 19 Feb 2025 16:59:27 +0100 Subject: [PATCH 2/4] Rework counting_iterator difference_type --- thrust/testing/counting_iterator.cu | 4 +- thrust/thrust/iterator/counting_iterator.h | 156 +++++---------------- 2 files changed, 39 insertions(+), 121 deletions(-) diff --git a/thrust/testing/counting_iterator.cu b/thrust/testing/counting_iterator.cu index 90abe1e4bc8..8a7a07bc7dd 100644 --- a/thrust/testing/counting_iterator.cu +++ b/thrust/testing/counting_iterator.cu @@ -23,8 +23,8 @@ static_assert(diff_type_is); static_assert(diff_type_is); static_assert(diff_type_is); #if _CCCL_HAS_INT128() -static_assert(diff_type_is<__int128_t, __int128_t>); -static_assert(diff_type_is<__uint128_t, long>); +static_assert(diff_type_is<__int128_t, ptrdiff_t>); +static_assert(diff_type_is<__uint128_t, ptrdiff_t>); #endif static_assert(diff_type_is); static_assert(diff_type_is); diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h index 4f159fa5432..6c448adaacb 100644 --- a/thrust/thrust/iterator/counting_iterator.h +++ b/thrust/thrust/iterator/counting_iterator.h @@ -46,69 +46,32 @@ #include #include -#include #include #include THRUST_NAMESPACE_BEGIN -// forward declaration of counting_iterator template class counting_iterator; namespace detail { -template -struct num_digits - : eval_if<::cuda::std::numeric_limits::is_specialized, - integral_constant::digits>, - integral_constant::digits - - (::cuda::std::numeric_limits::is_signed ? 1 : 0)>>::type -{}; // end num_digits - -template -struct integer_difference -//: eval_if< -// sizeof(Integer) >= sizeof(intmax_t), -// eval_if< -// is_signed::value, -// identity_, -// identity_ -// >, -// eval_if< -// sizeof(Integer) < sizeof(std::ptrdiff_t), -// identity_, -// identity_ -// > -// > -{ -private: - -public: - using type = - typename eval_if<::cuda::std::numeric_limits::is_signed - && (!::cuda::std::numeric_limits::is_bounded - || (int(::cuda::std::numeric_limits::digits) + 1 >= num_digits::value)), - identity_, - eval_if::digits) + 1 < num_digits::value, - identity_, - eval_if::digits) + 1 < num_digits::value, - identity_, - identity_>>>::type; -}; // end integer_difference - -template -struct numeric_difference - : eval_if<::cuda::std::is_integral::value, integer_difference, identity_> -{}; // end numeric_difference - template -_CCCL_HOST_DEVICE typename numeric_difference::type numeric_distance(Number x, Number y) -{ - using difference_type = typename numeric_difference::type; - return difference_type(y) - difference_type(x); -} // end numeric_distance +using counting_iterator_difference_type = ::cuda::std::_If< + ::cuda::std::is_integral_v, + // the difference between two int values can be larger than what an int can represent + ::cuda::std::_If, + // floating points use ptrdiff_t + ::cuda::std::_If<::cuda::std::is_floating_point_v, + ::cuda::std::ptrdiff_t, + // any other type, if it can represent the difference, can be used as difference type, + // otherwise also ptrdiff_t + ::cuda::std::_If<::cuda::std::numeric_limits::is_signed + && (!::cuda::std::numeric_limits::is_bounded + || ::cuda::std::numeric_limits::digits + > ::cuda::std::numeric_limits<::cuda::std::ptrdiff_t>::digits), + Number, + ::cuda::std::ptrdiff_t>>>; template struct make_counting_iterator_base @@ -116,19 +79,9 @@ struct make_counting_iterator_base using system = typename eval_if<::cuda::std::is_same::value, identity_, identity_>::type; - using traversal = replace_if_use_default< - Traversal, - eval_if::value, identity_, iterator_traversal>>; - - // unlike Boost, we explicitly use std::ptrdiff_t as the difference type - // for floating point counting_iterators + using traversal = replace_if_use_default>; using difference = - replace_if_use_default::value, - eval_if<::cuda::std::is_integral::value, - numeric_difference, - identity_<::cuda::std::ptrdiff_t>>, - lazy_trait>>; + typename replace_if_use_default>>::type; // our implementation departs from Boost's in that counting_iterator::dereference // returns a copy of its counter, rather than a reference to it. returning a reference @@ -142,50 +95,7 @@ struct make_counting_iterator_base traversal, Incrementable, difference>; -}; // end counting_iterator_base - -template -struct iterator_distance -{ - _CCCL_HOST_DEVICE static Difference distance(Incrementable1 x, Incrementable2 y) - { - return y - x; - } -}; - -template -struct number_distance -{ - _CCCL_HOST_DEVICE static Difference distance(Incrementable1 x, Incrementable2 y) - { - return static_cast(numeric_distance(x, y)); - } -}; - -template -struct counting_iterator_equal -{ - _CCCL_HOST_DEVICE static bool equal(Incrementable1 x, Incrementable2 y) - { - return x == y; - } -}; - -// specialization for floating point equality -template -struct counting_iterator_equal::value - || ::cuda::std::is_floating_point::value>> -{ - _CCCL_HOST_DEVICE static bool equal(Incrementable1 x, Incrementable2 y) - { - using d = number_distance; - return d::distance(x, y) == 0; - } }; - } // namespace detail //! \addtogroup iterators @@ -316,24 +226,32 @@ class _CCCL_DECLSPEC_EMPTY_BASES counting_iterator } // note that we implement equal specially for floating point counting_iterator - template + template _CCCL_HOST_DEVICE bool - equal(counting_iterator const& y) const + equal(counting_iterator const& y) const { - using e = detail::counting_iterator_equal; - return e::equal(this->base(), y.base()); + if constexpr (::cuda::is_floating_point_v) + { + return distance_to(y) == 0; + } + else + { + return this->base() == y.base(); + } } - template + template _CCCL_HOST_DEVICE difference_type - distance_to(counting_iterator const& y) const + distance_to(counting_iterator const& y) const { - using d = typename detail::eval_if< - detail::is_numeric::value, - detail::identity_>, - detail::identity_>>::type; - - return d::distance(this->base(), y.base()); + if constexpr (::cuda::std::is_integral::value) + { + return static_cast(y.base()) - static_cast(this->base()); + } + else + { + return y.base() - this->base(); + } } //! \endcond From 3ce4241f1e41e8c38763707167b5afa1a1bd4e9c Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 5 Mar 2025 18:59:02 +0100 Subject: [PATCH 3/4] Fix --- thrust/thrust/iterator/counting_iterator.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h index 6c448adaacb..9eb8bf9ee34 100644 --- a/thrust/thrust/iterator/counting_iterator.h +++ b/thrust/thrust/iterator/counting_iterator.h @@ -81,7 +81,7 @@ struct make_counting_iterator_base using traversal = replace_if_use_default>; using difference = - typename replace_if_use_default>>::type; + replace_if_use_default>>; // our implementation departs from Boost's in that counting_iterator::dereference // returns a copy of its counter, rather than a reference to it. returning a reference From afbb5d1f6198196938878b511cc7433a79c89ec5 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 5 Mar 2025 19:03:23 +0100 Subject: [PATCH 4/4] Simplify --- thrust/thrust/iterator/counting_iterator.h | 17 ++--------------- 1 file changed, 2 insertions(+), 15 deletions(-) diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h index 9eb8bf9ee34..f06499b9d37 100644 --- a/thrust/thrust/iterator/counting_iterator.h +++ b/thrust/thrust/iterator/counting_iterator.h @@ -57,21 +57,8 @@ class counting_iterator; namespace detail { template -using counting_iterator_difference_type = ::cuda::std::_If< - ::cuda::std::is_integral_v, - // the difference between two int values can be larger than what an int can represent - ::cuda::std::_If, - // floating points use ptrdiff_t - ::cuda::std::_If<::cuda::std::is_floating_point_v, - ::cuda::std::ptrdiff_t, - // any other type, if it can represent the difference, can be used as difference type, - // otherwise also ptrdiff_t - ::cuda::std::_If<::cuda::std::numeric_limits::is_signed - && (!::cuda::std::numeric_limits::is_bounded - || ::cuda::std::numeric_limits::digits - > ::cuda::std::numeric_limits<::cuda::std::ptrdiff_t>::digits), - Number, - ::cuda::std::ptrdiff_t>>>; +using counting_iterator_difference_type = + ::cuda::std::_If<::cuda::std::is_integral_v && sizeof(Number) < sizeof(int), int, ::cuda::std::ptrdiff_t>; template struct make_counting_iterator_base