diff --git a/thrust/testing/counting_iterator.cu b/thrust/testing/counting_iterator.cu index a7fca340523..8a7a07bc7dd 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, ptrdiff_t>); +static_assert(diff_type_is<__uint128_t, ptrdiff_t>); +#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 diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h index 4f159fa5432..f06499b9d37 100644 --- a/thrust/thrust/iterator/counting_iterator.h +++ b/thrust/thrust/iterator/counting_iterator.h @@ -46,69 +46,19 @@ #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 && sizeof(Number) < sizeof(int), int, ::cuda::std::ptrdiff_t>; template struct make_counting_iterator_base @@ -116,19 +66,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>>; + 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 @@ -142,50 +82,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 +213,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