From 8e16c1067c251efe0e8e24deab251276ade5a934 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 5 Mar 2025 17:49:47 +0100 Subject: [PATCH] refactor --- thrust/testing/counting_iterator.cu | 15 -------- thrust/thrust/iterator/counting_iterator.h | 44 +++++++++++++--------- 2 files changed, 27 insertions(+), 32 deletions(-) diff --git a/thrust/testing/counting_iterator.cu b/thrust/testing/counting_iterator.cu index 372cbcd5995..e1c281d3c4c 100644 --- a/thrust/testing/counting_iterator.cu +++ b/thrust/testing/counting_iterator.cu @@ -321,19 +321,4 @@ void TestCountingIteratorPointer() } DECLARE_UNITTEST(TestCountingIteratorPointer); -void TestCountingIteratorStridedPointer() -{ - ::cuda::std::array, 4> arr{{{1, 2}, {3, 4}, {5, 6}, {7, 8}}}; - - // iterate over all second elements - auto iter = thrust::make_counting_iterator( - &arr[0].second, ::cuda::std::integral_constant)>{}); - - cuda::std::fill(iter, iter + 4, 1337); - - const ::cuda::std::array, 4> reference{{{1, 1337}, {3, 1337}, {5, 1337}, {7, 1337}}}; - ASSERT_EQUAL(arr, reference); -} -DECLARE_UNITTEST(TestCountingIteratorStridedPointer); - _CCCL_DIAG_POP diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h index b7c2a6a6016..725039916c2 100644 --- a/thrust/thrust/iterator/counting_iterator.h +++ b/thrust/thrust/iterator/counting_iterator.h @@ -188,17 +188,6 @@ struct counting_iterator_equal -struct value_holder -{ - T value; - - _CCCL_HOST_DEVICE auto operator()() const - { - return value; - } -}; } // namespace detail //! \addtogroup iterators @@ -336,7 +325,7 @@ class _CCCL_DECLSPEC_EMPTY_BASES counting_iterator auto step() const { static_assert(!::cuda::std::is_same_v); - return static_cast(*this)(); + return static_cast(*this).value; } _CCCL_EXEC_CHECK_DISABLE @@ -422,25 +411,46 @@ inline _CCCL_HOST_DEVICE counting_iterator make_counting_iterator return counting_iterator(x); } +namespace detail +{ +// Holds a runtime step +template +struct runtime_step_holder +{ + T value; +}; + +// Holds a compile-time step +// (we cannot use ::cuda::std::integral_constant, because it has a conversion operator to T that causes an ambiguity +// with operator+(counting_iterator, counting_iterator::difference_type) in any expression `counting_iterator + +// integral`. +template +struct compile_time_step_holder +{ + static constexpr T value = Value; +}; +} // namespace detail + +//! Constructs a counting_iterator with a runtime stride template -inline _CCCL_HOST_DEVICE auto make_counting_iterator(Incrementable x, Stride stride) +_CCCL_HOST_DEVICE auto make_counting_iterator(Incrementable x, Stride stride) { return counting_iterator>(x, {stride}); + detail::runtime_step_holder>(x, {stride}); } +//! Constructs a counting_iterator with a compile-time stride template -inline _CCCL_HOST_DEVICE auto -make_counting_iterator(Incrementable x, ::cuda::std::integral_constant stride) +_CCCL_HOST_DEVICE auto make_counting_iterator(Incrementable x, ::cuda::std::integral_constant) { return counting_iterator>(x, stride); + detail::compile_time_step_holder>(x, {}); } //! \} // end fancyiterators