Skip to content

Inherit thrust::transform_iterator traversal from base iterator traversal #5883

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

Merged
bernhardmgruber merged 3 commits into from
Sep 16, 2025
Merged

Inherit thrust::transform_iterator traversal from base iterator traversal #5883

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
56c7179
Inherit thrust::transform_iterator traversal from base iterator trave…
bernhardmgruber Sep 15, 2025
763c837
More tests
bernhardmgruber Sep 15, 2025
1ed55d1
Also test traversal
bernhardmgruber Sep 15, 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
159 changes: 124 additions & 35 deletions thrust/testing/catch2_test_iterator_traits.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,8 @@
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>

#include <vector>

#include "catch2_test_helper.h"

template <typename Cat, typename Tag>
Expand Down Expand Up @@ -43,53 +45,140 @@ TEST_CASE("iterator_category_to_traversal", "[iterators]")
STATIC_REQUIRE(traversal_is<thrust::random_access_host_iterator_tag, thrust::random_access_traversal_tag>);
}

TEST_CASE("iterator system propagation - any system", "[iterators]")
struct cuda_make_counting_iterator
{
[[maybe_unused]] auto counting_it = thrust::make_counting_iterator(0);
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(counting_it)>, thrust::any_system_tag>);

[[maybe_unused]] auto transform_it = thrust::make_transform_iterator(counting_it, thrust::square<>{});
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(transform_it)>, thrust::any_system_tag>);

[[maybe_unused]] auto zip_it = thrust::make_zip_iterator(counting_it, transform_it);
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(zip_it)>, thrust::any_system_tag>);
template <typename... Args>
auto operator()(Args&&... args) const
{
return cuda::make_counting_iterator(cuda::std::forward<Args>(args)...);
}
};

struct cuda_make_transform_iterator
{
template <typename... Args>
auto operator()(Args&&... args) const
{
return cuda::make_transform_iterator(cuda::std::forward<Args>(args)...);
}
};

struct cuda_make_zip_iterator
{
template <typename... Args>
auto operator()(Args&&... args) const
{
return cuda::make_zip_iterator(cuda::std::forward<Args>(args)...);
}
};

struct thrust_make_counting_iterator
{
template <typename... Args>
auto operator()(Args&&... args) const
{
return thrust::make_counting_iterator(cuda::std::forward<Args>(args)...);
}
};

struct thrust_make_transform_iterator
{
template <typename... Args>
auto operator()(Args&&... args) const
{
return thrust::make_transform_iterator(cuda::std::forward<Args>(args)...);
}
};

struct thrust_make_zip_iterator
{
template <typename... Args>
auto operator()(Args&&... args) const
{
return thrust::make_zip_iterator(cuda::std::forward<Args>(args)...);
}
};

template <typename Iterator>
inline constexpr bool has_random_access_traversal =
cuda::std::is_same_v<thrust::iterator_traversal_t<Iterator>, thrust::random_access_traversal_tag>;

using cuda::std::__type_cartesian_product;
using cuda::std::__type_list;

using make_counting_t = __type_list<cuda_make_counting_iterator, thrust_make_counting_iterator>;
using make_transform_t = __type_list<cuda_make_transform_iterator, thrust_make_transform_iterator>;
using make_zip_t = __type_list<cuda_make_zip_iterator, thrust_make_zip_iterator>;
using make_it_t = __type_cartesian_product<make_counting_t, make_transform_t, make_zip_t>;
TEMPLATE_LIST_TEST_CASE("iterator system and traversal propagation - any system", "[iterators]", make_it_t)
{
using namespace cuda::std;
__type_at_c<0, TestType> make_counting_iterator;
__type_at_c<1, TestType> make_transform_iterator;
__type_at_c<2, TestType> make_zip_iterator;

auto counting_it = make_counting_iterator(0);
STATIC_REQUIRE(is_same_v<thrust::iterator_system_t<decltype(counting_it)>, thrust::any_system_tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(counting_it)>);

auto transform_it = make_transform_iterator(counting_it, thrust::square<>{});
STATIC_REQUIRE(is_same_v<thrust::iterator_system_t<decltype(transform_it)>, thrust::any_system_tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(transform_it)>);

[[maybe_unused]] auto zip_it = make_zip_iterator(counting_it, transform_it);
STATIC_REQUIRE(is_same_v<thrust::iterator_system_t<decltype(zip_it)>, thrust::any_system_tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(zip_it)>);
}

TEST_CASE("iterator system propagation - host system", "[iterators]")
{
[[maybe_unused]] auto h_vec = thrust::host_vector<int>{};
[[maybe_unused]] auto h_vec_it = h_vec.begin();
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(h_vec_it)>, thrust::host_system_tag>);
auto expected_tag(thrust::device_vector<int>) -> thrust::device_system_tag;
auto expected_tag(thrust::host_vector<int>) -> thrust::host_system_tag;
auto expected_tag(std::vector<int>) -> thrust::host_system_tag;

[[maybe_unused]] auto transform_it = thrust::make_transform_iterator(h_vec_it, thrust::square<>{});
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(transform_it)>, thrust::host_system_tag>);
using vectors = __type_list<thrust::device_vector<int>, thrust::host_vector<int>, std::vector<int>>;
using make_vec_and_it_t = __type_cartesian_product<vectors, make_transform_t, make_zip_t>;
TEMPLATE_LIST_TEST_CASE("iterator system and traversal propagation - vectors", "[iterators]", make_vec_and_it_t)
{
using namespace cuda::std;
__type_at_c<0, TestType> vec{};
__type_at_c<1, TestType> make_transform_iterator;
__type_at_c<2, TestType> make_zip_iterator;

[[maybe_unused]] auto zip_it = thrust::make_zip_iterator(h_vec_it, transform_it);
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(zip_it)>, thrust::host_system_tag>);
}
using tag = decltype(expected_tag(vec));

TEST_CASE("iterator system propagation - device system", "[iterators]")
{
[[maybe_unused]] auto d_vec = thrust::device_vector<int>{};
[[maybe_unused]] auto d_vec_it = d_vec.begin();
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(d_vec_it)>, thrust::device_system_tag>);
auto vec_it = vec.begin();
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(vec_it)>, tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(vec_it)>);

[[maybe_unused]] auto transform_it = thrust::make_transform_iterator(d_vec_it, thrust::square<>{});
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(transform_it)>, thrust::device_system_tag>);
using thrust::placeholders::_1;
auto transform_it = make_transform_iterator(vec_it, _1 + 1);
STATIC_REQUIRE(is_same_v<thrust::iterator_system_t<decltype(transform_it)>, tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(transform_it)>);

[[maybe_unused]] auto zip_it = thrust::make_zip_iterator(d_vec_it, transform_it);
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(zip_it)>, thrust::device_system_tag>);
[[maybe_unused]] auto zip_it = make_zip_iterator(vec_it, transform_it);
STATIC_REQUIRE(is_same_v<thrust::iterator_system_t<decltype(zip_it)>, tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(zip_it)>);
}

TEST_CASE("iterator system propagation - any and device system", "[iterators]")
using make_vec_count_zip_it_t = __type_cartesian_product<vectors, make_counting_t, make_zip_t>;
TEMPLATE_LIST_TEST_CASE(
"iterator system and traversal propagation - vectors and any system", "[iterators]", make_vec_count_zip_it_t)
{
[[maybe_unused]] auto d_vec = thrust::device_vector<int>{};
[[maybe_unused]] auto d_vec_it = d_vec.begin();
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(d_vec_it)>, thrust::device_system_tag>);
using namespace cuda::std;
__type_at_c<0, TestType> vec{};
__type_at_c<1, TestType> make_counting_iterator;
__type_at_c<2, TestType> make_zip_iterator;

using vec_tag = decltype(expected_tag(vec));

auto vec_it = vec.begin();
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(vec_it)>, vec_tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(vec_it)>);

[[maybe_unused]] auto counting_it = thrust::make_counting_iterator(0);
auto counting_it = make_counting_iterator(0);
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(counting_it)>, thrust::any_system_tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(counting_it)>);

[[maybe_unused]] auto zip_it = thrust::make_zip_iterator(d_vec_it, counting_it);
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(zip_it)>, thrust::device_system_tag>);
[[maybe_unused]] auto zip_it = make_zip_iterator(vec_it, counting_it);
STATIC_REQUIRE(cuda::std::is_same_v<thrust::iterator_system_t<decltype(zip_it)>, vec_tag>);
STATIC_REQUIRE(has_random_access_traversal<decltype(zip_it)>);
}
4 changes: 2 additions & 2 deletions thrust/thrust/iterator/transform_iterator.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -93,8 +93,8 @@ struct make_transform_iterator_base
iterator_adaptor<transform_iterator<UnaryFunc, Iterator, Reference, Value>,
Iterator,
value_type,
use_default,
typename ::cuda::std::iterator_traits<Iterator>::iterator_category,
use_default, // pick system from Iterator
use_default, // pick traversal from Iterator
Copy link
Contributor Author

@bernhardmgruber bernhardmgruber Sep 15, 2025

Choose a reason for hiding this comment

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

This line is the fix to #5860 and I also think it's an overall cleaner behavior. I don't know though why the old behavior was chosen. Need to do some digging.

Copy link
Contributor Author

@bernhardmgruber bernhardmgruber Sep 15, 2025

Choose a reason for hiding this comment

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

Found a hint: https://github.com/NVIDIA/cccl/pull/1915/files#diff-866ad8d2535a5dd3e954db56809b4673fcce2b064fc630fe24c872c5e8195a81L61-L67

  typedef thrust::iterator_adaptor<transform_iterator<UnaryFunc, Iterator, Reference, Value>,
                                   Iterator,
                                   value_type,
                                   thrust::use_default // Leave the system alone
                                                       //, thrust::use_default   // Leave the traversal alone
                                                       // use the Iterator's category to let any system iterators remain
                                                       // random access even though transform_iterator's reference type
                                                       // may not be a reference
                                                       // XXX figure out why only iterators whose reference types are
                                                       // true references are random access
                                   ,
                                   typename thrust::iterator_traits<Iterator>::iterator_category,
                                   reference>

It seems category was used to work around any system iterators loosing their random accessness. Let's see whether this problem still exists.

Copy link
Contributor Author

@bernhardmgruber bernhardmgruber Sep 15, 2025

Choose a reason for hiding this comment

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

I added tests to check whether the properties of a random access any system iterator are retained through a transform iterator. And they are. So I think the fix is legit.

reference>;
};

Expand Down