Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Rework counting_iterator difference #3861

Merged
merged 4 commits into from
Mar 6, 2025
Merged

Rework counting_iterator difference #3861

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
b9dea3d
Add tests for counting_iterator difference_type
bernhardmgruber Feb 19, 2025
22d84f4
Rework counting_iterator difference_type
bernhardmgruber Feb 19, 2025
3ce4241
Fix
bernhardmgruber Mar 5, 2025
afbb5d1
Simplify
bernhardmgruber Mar 5, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
28 changes: 28 additions & 0 deletions thrust/testing/counting_iterator.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,34 @@

#include <unittest/unittest.h>

template <typename ValueType, typename DifferenceType>
inline constexpr bool diff_type_is =
::cuda::std::is_same_v<typename thrust::counting_iterator<ValueType>::difference_type, DifferenceType>;

static_assert(diff_type_is<int8_t, int>);
static_assert(diff_type_is<uint8_t, int>);
static_assert(diff_type_is<int16_t, int>);
static_assert(diff_type_is<uint16_t, int>);
static_assert(diff_type_is<int32_t, ptrdiff_t>);
static_assert(diff_type_is<uint32_t, ptrdiff_t>);
static_assert(diff_type_is<int64_t, ptrdiff_t>);
static_assert(diff_type_is<uint64_t, ptrdiff_t>);
#if _CCCL_HAS_INT128()
static_assert(diff_type_is<__int128_t, ptrdiff_t>);
static_assert(diff_type_is<__uint128_t, ptrdiff_t>);
Comment on lines +26 to +27
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Before the second commit of this PR (the first commit adds this test), the answer was:

static_assert(diff_type_is<__int128_t, __int128_t>);
static_assert(diff_type_is<__uint128_t, long>);

Which may be surprising.

#endif
static_assert(diff_type_is<float, ptrdiff_t>);
static_assert(diff_type_is<double, ptrdiff_t>);

struct custom_int
{
_CCCL_HOST_DEVICE custom_int(int) {}
_CCCL_HOST_DEVICE operator int() const;
};
static_assert(thrust::detail::is_numeric<custom_int>::value);

static_assert(diff_type_is<custom_int, ptrdiff_t>);

_CCCL_DIAG_PUSH
_CCCL_DIAG_SUPPRESS_MSVC(4244 4267) // possible loss of data

Expand Down
143 changes: 24 additions & 119 deletions thrust/thrust/iterator/counting_iterator.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,89 +46,29 @@
#include <thrust/iterator/iterator_traits.h>

#include <cuda/std/cstddef>
#include <cuda/std/limits>
#include <cuda/std/type_traits>
#include <cuda/type_traits>

THRUST_NAMESPACE_BEGIN

// forward declaration of counting_iterator
template <typename Incrementable, typename System, typename Traversal, typename Difference>
class counting_iterator;

namespace detail
{
template <typename T>
struct num_digits
: eval_if<::cuda::std::numeric_limits<T>::is_specialized,
integral_constant<int, ::cuda::std::numeric_limits<T>::digits>,
integral_constant<int,
sizeof(T) * ::cuda::std::numeric_limits<unsigned char>::digits
- (::cuda::std::numeric_limits<T>::is_signed ? 1 : 0)>>::type
{}; // end num_digits

template <typename Integer>
struct integer_difference
//: eval_if<
// sizeof(Integer) >= sizeof(intmax_t),
// eval_if<
// is_signed<Integer>::value,
// identity_<Integer>,
// identity_<intmax_t>
// >,
// eval_if<
// sizeof(Integer) < sizeof(std::ptrdiff_t),
// identity_<std::ptrdiff_t>,
// identity_<intmax_t>
// >
// >
{
private:

public:
using type =
typename eval_if<::cuda::std::numeric_limits<Integer>::is_signed
&& (!::cuda::std::numeric_limits<Integer>::is_bounded
|| (int(::cuda::std::numeric_limits<Integer>::digits) + 1 >= num_digits<intmax_t>::value)),
identity_<Integer>,
eval_if<int(::cuda::std::numeric_limits<Integer>::digits) + 1 < num_digits<int>::value,
identity_<int>,
eval_if<int(::cuda::std::numeric_limits<Integer>::digits) + 1 < num_digits<long>::value,
identity_<long>,
identity_<intmax_t>>>>::type;
}; // end integer_difference

template <typename Number>
struct numeric_difference
: eval_if<::cuda::std::is_integral<Number>::value, integer_difference<Number>, identity_<Number>>
{}; // end numeric_difference

template <typename Number>
_CCCL_HOST_DEVICE typename numeric_difference<Number>::type numeric_distance(Number x, Number y)
{
using difference_type = typename numeric_difference<Number>::type;
return difference_type(y) - difference_type(x);
} // end numeric_distance
using counting_iterator_difference_type =
::cuda::std::_If<::cuda::std::is_integral_v<Number> && sizeof(Number) < sizeof(int), int, ::cuda::std::ptrdiff_t>;

template <typename Incrementable, typename System, typename Traversal, typename Difference>
struct make_counting_iterator_base
{
using system =
typename eval_if<::cuda::std::is_same<System, use_default>::value, identity_<any_system_tag>, identity_<System>>::type;

using traversal = replace_if_use_default<
Traversal,
eval_if<is_numeric<Incrementable>::value, identity_<random_access_traversal_tag>, iterator_traversal<Incrementable>>>;

// unlike Boost, we explicitly use std::ptrdiff_t as the difference type
// for floating point counting_iterators
using traversal = replace_if_use_default<Traversal, ::cuda::std::type_identity<random_access_traversal_tag>>;
using difference =
replace_if_use_default<Difference,
eval_if<is_numeric<Incrementable>::value,
eval_if<::cuda::std::is_integral<Incrementable>::value,
numeric_difference<Incrementable>,
identity_<::cuda::std::ptrdiff_t>>,
lazy_trait<it_difference_t, Incrementable>>>;
replace_if_use_default<Difference, ::cuda::std::type_identity<counting_iterator_difference_type<Incrementable>>>;

// 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
Expand All @@ -142,50 +82,7 @@ struct make_counting_iterator_base
traversal,
Incrementable,
difference>;
}; // end counting_iterator_base

template <typename Difference, typename Incrementable1, typename Incrementable2>
struct iterator_distance
{
_CCCL_HOST_DEVICE static Difference distance(Incrementable1 x, Incrementable2 y)
{
return y - x;
}
};

template <typename Difference, typename Incrementable1, typename Incrementable2>
struct number_distance
{
_CCCL_HOST_DEVICE static Difference distance(Incrementable1 x, Incrementable2 y)
{
return static_cast<Difference>(numeric_distance(x, y));
}
};

template <typename Difference, typename Incrementable1, typename Incrementable2, typename Enable = void>
struct counting_iterator_equal
{
_CCCL_HOST_DEVICE static bool equal(Incrementable1 x, Incrementable2 y)
{
return x == y;
}
};

// specialization for floating point equality
template <typename Difference, typename Incrementable1, typename Incrementable2>
struct counting_iterator_equal<Difference,
Incrementable1,
Incrementable2,
::cuda::std::enable_if_t<::cuda::std::is_floating_point<Incrementable1>::value
|| ::cuda::std::is_floating_point<Incrementable2>::value>>
{
_CCCL_HOST_DEVICE static bool equal(Incrementable1 x, Incrementable2 y)
{
using d = number_distance<Difference, Incrementable1, Incrementable2>;
return d::distance(x, y) == 0;
}
};

} // namespace detail

//! \addtogroup iterators
Expand Down Expand Up @@ -316,24 +213,32 @@ class _CCCL_DECLSPEC_EMPTY_BASES counting_iterator
}

// note that we implement equal specially for floating point counting_iterator
template <typename OtherIncrementable, typename OtherSystem, typename OtherTraversal, typename OtherDifference>
template <typename OtherSystem, typename OtherTraversal, typename OtherDifference>
_CCCL_HOST_DEVICE bool
equal(counting_iterator<OtherIncrementable, OtherSystem, OtherTraversal, OtherDifference> const& y) const
equal(counting_iterator<Incrementable, OtherSystem, OtherTraversal, OtherDifference> const& y) const
{
using e = detail::counting_iterator_equal<difference_type, Incrementable, OtherIncrementable>;
return e::equal(this->base(), y.base());
if constexpr (::cuda::is_floating_point_v<Incrementable>)
{
return distance_to(y) == 0;
}
else
{
return this->base() == y.base();
}
}

template <class OtherIncrementable>
template <typename OtherSystem, typename OtherTraversal, typename OtherDifference>
_CCCL_HOST_DEVICE difference_type
distance_to(counting_iterator<OtherIncrementable, System, Traversal, Difference> const& y) const
distance_to(counting_iterator<Incrementable, OtherSystem, OtherTraversal, OtherDifference> const& y) const
{
using d = typename detail::eval_if<
detail::is_numeric<Incrementable>::value,
detail::identity_<detail::number_distance<difference_type, Incrementable, OtherIncrementable>>,
detail::identity_<detail::iterator_distance<difference_type, Incrementable, OtherIncrementable>>>::type;

return d::distance(this->base(), y.base());
if constexpr (::cuda::std::is_integral<Incrementable>::value)
{
return static_cast<difference_type>(y.base()) - static_cast<difference_type>(this->base());
}
else
{
return y.base() - this->base();
}
}

//! \endcond
Expand Down
Loading