diff --git a/libcudacxx/include/cuda/std/__algorithm/equal.h b/libcudacxx/include/cuda/std/__algorithm/equal.h index a7b427e431a..452a198b286 100644 --- a/libcudacxx/include/cuda/std/__algorithm/equal.h +++ b/libcudacxx/include/cuda/std/__algorithm/equal.h @@ -103,8 +103,8 @@ equal(_InputIterator1 __first1, __first2, __last2, __pred, - typename iterator_traits<_InputIterator1>::iterator_category(), - typename iterator_traits<_InputIterator2>::iterator_category()); + __iterator_traits_category_or_concept_t<_InputIterator1>(), + __iterator_traits_category_or_concept_t<_InputIterator2>()); } template @@ -117,8 +117,8 @@ equal(_InputIterator1 __first1, _InputIterator1 __last1, _InputIterator2 __first __first2, __last2, __equal_to{}, - typename iterator_traits<_InputIterator1>::iterator_category(), - typename iterator_traits<_InputIterator2>::iterator_category()); + __iterator_traits_category_or_concept_t<_InputIterator1>(), + __iterator_traits_category_or_concept_t<_InputIterator2>()); } _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__algorithm/fill.h b/libcudacxx/include/cuda/std/__algorithm/fill.h index 113366a9017..6bc1882b2f3 100644 --- a/libcudacxx/include/cuda/std/__algorithm/fill.h +++ b/libcudacxx/include/cuda/std/__algorithm/fill.h @@ -48,7 +48,7 @@ __fill(_RandomAccessIterator __first, _RandomAccessIterator __last, const _Tp& _ template _CCCL_API constexpr void fill(_ForwardIterator __first, _ForwardIterator __last, const _Tp& __value_) { - ::cuda::std::__fill(__first, __last, __value_, typename iterator_traits<_ForwardIterator>::iterator_category()); + ::cuda::std::__fill(__first, __last, __value_, __iterator_traits_category_or_concept_t<_ForwardIterator>()); } _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__algorithm/find_end.h b/libcudacxx/include/cuda/std/__algorithm/find_end.h index 80f7ab32d48..1f6dfaecfdc 100644 --- a/libcudacxx/include/cuda/std/__algorithm/find_end.h +++ b/libcudacxx/include/cuda/std/__algorithm/find_end.h @@ -207,8 +207,8 @@ template ::iterator_category{}, - typename iterator_traits<_ForwardIterator2>::iterator_category{}); + __iterator_traits_category_or_concept_t<_ForwardIterator1>{}, + __iterator_traits_category_or_concept_t<_ForwardIterator2>{}); } template diff --git a/libcudacxx/include/cuda/std/__algorithm/iterator_operations.h b/libcudacxx/include/cuda/std/__algorithm/iterator_operations.h index de305d02fdd..562bb7e314f 100644 --- a/libcudacxx/include/cuda/std/__algorithm/iterator_operations.h +++ b/libcudacxx/include/cuda/std/__algorithm/iterator_operations.h @@ -55,9 +55,6 @@ struct _IterOps<_RangeAlgPolicy> template using __value_type = iter_value_t<_Iter>; - template - using __iterator_category = ::cuda::std::ranges::__iterator_concept<_Iter>; - template using __difference_type = iter_difference_t<_Iter>; @@ -79,9 +76,6 @@ struct _IterOps<_ClassicAlgPolicy> template using __value_type = typename iterator_traits<_Iter>::value_type; - template - using __iterator_category = typename iterator_traits<_Iter>::iterator_category; - template using __difference_type = typename iterator_traits<_Iter>::difference_type; diff --git a/libcudacxx/include/cuda/std/__algorithm/partition.h b/libcudacxx/include/cuda/std/__algorithm/partition.h index 017d271e520..6dae91a5b20 100644 --- a/libcudacxx/include/cuda/std/__algorithm/partition.h +++ b/libcudacxx/include/cuda/std/__algorithm/partition.h @@ -107,7 +107,7 @@ _CCCL_EXEC_CHECK_DISABLE template _CCCL_API constexpr _ForwardIterator partition(_ForwardIterator __first, _ForwardIterator __last, _Predicate __pred) { - using _IterCategory = typename iterator_traits<_ForwardIterator>::iterator_category; + using _IterCategory = __iterator_traits_category_or_concept_t<_ForwardIterator>; auto __result = ::cuda::std::__partition<_ClassicAlgPolicy>( ::cuda::std::move(__first), ::cuda::std::move(__last), __pred, _IterCategory()); return __result.first; diff --git a/libcudacxx/include/cuda/std/__algorithm/reverse.h b/libcudacxx/include/cuda/std/__algorithm/reverse.h index 93bf061db58..b11658568e5 100644 --- a/libcudacxx/include/cuda/std/__algorithm/reverse.h +++ b/libcudacxx/include/cuda/std/__algorithm/reverse.h @@ -63,7 +63,7 @@ _CCCL_EXEC_CHECK_DISABLE template _CCCL_API constexpr void __reverse(_BidirectionalIterator __first, _Sentinel __last) { - using _IterCategory = typename _IterOps<_AlgPolicy>::template __iterator_category<_BidirectionalIterator>; + using _IterCategory = __iterator_traits_category_or_concept_t<_BidirectionalIterator>; ::cuda::std::__reverse_impl<_AlgPolicy>(::cuda::std::move(__first), ::cuda::std::move(__last), _IterCategory()); } diff --git a/libcudacxx/include/cuda/std/__algorithm/rotate.h b/libcudacxx/include/cuda/std/__algorithm/rotate.h index e894e3b1ce4..640e869187c 100644 --- a/libcudacxx/include/cuda/std/__algorithm/rotate.h +++ b/libcudacxx/include/cuda/std/__algorithm/rotate.h @@ -238,7 +238,7 @@ _CCCL_API constexpr pair<_Iterator, _Iterator> __rotate(_Iterator __first, _Iter return _Ret(::cuda::std::move(__first), ::cuda::std::move(__last_iter)); } - using _IterCategory = typename _IterOps<_AlgPolicy>::template __iterator_category<_Iterator>; + using _IterCategory = __iterator_traits_category_or_concept_t<_Iterator>; auto __result = ::cuda::std::__rotate_impl<_AlgPolicy>( ::cuda::std::move(__first), ::cuda::std::move(__middle), __last_iter, _IterCategory()); diff --git a/libcudacxx/include/cuda/std/__algorithm/sample.h b/libcudacxx/include/cuda/std/__algorithm/sample.h index 32370b77ccb..418522c6566 100644 --- a/libcudacxx/include/cuda/std/__algorithm/sample.h +++ b/libcudacxx/include/cuda/std/__algorithm/sample.h @@ -89,7 +89,7 @@ _CCCL_API _SampleIterator __sample( _Distance __n, _UniformRandomNumberGenerator& __g) { - using _PopCategory = typename iterator_traits<_PopulationIterator>::iterator_category; + using _PopCategory = __iterator_traits_category_or_concept_t<_PopulationIterator>; using _Difference = typename iterator_traits<_PopulationIterator>::difference_type; static_assert(__has_forward_traversal<_PopulationIterator> || __has_random_access_traversal<_SampleIterator>, "SampleIterator must meet the requirements of RandomAccessIterator"); diff --git a/libcudacxx/include/cuda/std/__algorithm/search.h b/libcudacxx/include/cuda/std/__algorithm/search.h index efd24317189..6c6e35f7940 100644 --- a/libcudacxx/include/cuda/std/__algorithm/search.h +++ b/libcudacxx/include/cuda/std/__algorithm/search.h @@ -159,8 +159,8 @@ search(_ForwardIterator1 __first1, __first2, __last2, __pred, - typename iterator_traits<_ForwardIterator1>::iterator_category(), - typename iterator_traits<_ForwardIterator2>::iterator_category()) + __iterator_traits_category_or_concept_t<_ForwardIterator1>(), + __iterator_traits_category_or_concept_t<_ForwardIterator2>()) .first; } diff --git a/libcudacxx/include/cuda/std/__algorithm/search_n.h b/libcudacxx/include/cuda/std/__algorithm/search_n.h index cba94f2ecac..5bdaecba3b7 100644 --- a/libcudacxx/include/cuda/std/__algorithm/search_n.h +++ b/libcudacxx/include/cuda/std/__algorithm/search_n.h @@ -146,7 +146,7 @@ search_n(_ForwardIterator __first, _ForwardIterator __last, _Size __count, const __convert_to_integral(__count), __value_, __pred, - typename iterator_traits<_ForwardIterator>::iterator_category()); + __iterator_traits_category_or_concept_t<_ForwardIterator>()); } template diff --git a/libcudacxx/include/cuda/std/__algorithm/shift_left.h b/libcudacxx/include/cuda/std/__algorithm/shift_left.h index 10c66009929..45769448513 100644 --- a/libcudacxx/include/cuda/std/__algorithm/shift_left.h +++ b/libcudacxx/include/cuda/std/__algorithm/shift_left.h @@ -73,7 +73,7 @@ _CCCL_API constexpr _ForwardIterator shift_left( return __last; } - using _IterCategory = typename iterator_traits<_ForwardIterator>::iterator_category; + using _IterCategory = __iterator_traits_category_or_concept_t<_ForwardIterator>; return __shift_left(__first, __last, __n, _IterCategory()); } diff --git a/libcudacxx/include/cuda/std/__algorithm/shift_right.h b/libcudacxx/include/cuda/std/__algorithm/shift_right.h index c6d38866049..167b43942dc 100644 --- a/libcudacxx/include/cuda/std/__algorithm/shift_right.h +++ b/libcudacxx/include/cuda/std/__algorithm/shift_right.h @@ -133,7 +133,7 @@ _CCCL_API constexpr _ForwardIterator shift_right( return __first; } - using _IterCategory = typename iterator_traits<_ForwardIterator>::iterator_category; + using _IterCategory = __iterator_traits_category_or_concept_t<_ForwardIterator>; return __shift_right(__first, __last, __n, _IterCategory()); } diff --git a/libcudacxx/include/cuda/std/__algorithm/stable_partition.h b/libcudacxx/include/cuda/std/__algorithm/stable_partition.h index 3656fa1bf40..1d2898febe3 100644 --- a/libcudacxx/include/cuda/std/__algorithm/stable_partition.h +++ b/libcudacxx/include/cuda/std/__algorithm/stable_partition.h @@ -347,7 +347,7 @@ _CCCL_API _ForwardIterator __stable_partition( template _CCCL_API _ForwardIterator stable_partition(_ForwardIterator __first, _ForwardIterator __last, _Predicate __pred) { - using _IterCategory = typename iterator_traits<_ForwardIterator>::iterator_category; + using _IterCategory = __iterator_traits_category_or_concept_t<_ForwardIterator>; return ::cuda::std::__stable_partition<_ClassicAlgPolicy, _Predicate&>( ::cuda::std::move(__first), ::cuda::std::move(__last), __pred, _IterCategory()); } diff --git a/libcudacxx/include/cuda/std/__functional/default_searcher.h b/libcudacxx/include/cuda/std/__functional/default_searcher.h index 9928540e343..9cc62430941 100644 --- a/libcudacxx/include/cuda/std/__functional/default_searcher.h +++ b/libcudacxx/include/cuda/std/__functional/default_searcher.h @@ -56,8 +56,8 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT default_searcher __first_, __last_, __pred_, - typename ::cuda::std::iterator_traits<_ForwardIterator>::iterator_category(), - typename ::cuda::std::iterator_traits<_ForwardIterator2>::iterator_category()); + __iterator_traits_category_or_concept_t<_ForwardIterator>(), + __iterator_traits_category_or_concept_t<_ForwardIterator2>()); } private: diff --git a/libcudacxx/include/cuda/std/__iterator/iterator_traits.h b/libcudacxx/include/cuda/std/__iterator/iterator_traits.h index f147f615263..974db1c41ab 100644 --- a/libcudacxx/include/cuda/std/__iterator/iterator_traits.h +++ b/libcudacxx/include/cuda/std/__iterator/iterator_traits.h @@ -238,6 +238,9 @@ _CCCL_CONCEPT __has_member_pointer = _CCCL_REQUIRES_EXPR((_Tp))(typename(typenam template _CCCL_CONCEPT __has_member_iterator_category = _CCCL_REQUIRES_EXPR((_Tp))(typename(typename _Tp::iterator_category)); +template +_CCCL_CONCEPT __has_member_iterator_concept = _CCCL_REQUIRES_EXPR((_Tp))(typename(typename _Tp::iterator_concept)); + // The `cpp17-*-iterator` exposition-only concepts have very similar names to the `Cpp17*Iterator` named requirements // from `[iterator.cpp17]`. To avoid confusion between the two, the exposition-only concepts have been banished to // a "detail" namespace indicating they have a niche use-case. @@ -569,6 +572,45 @@ using __iter_diff_t = typename iterator_traits<_Iter>::difference_type; template using __iter_value_type = typename iterator_traits<_Iter>::value_type; +// C++20 iterators do not play nicely with C++17 interfaces, because they commmonly use `iterator_concept` to +// communicate their iterator category. Deduce the actual iterator category from the maximum of `iterator_concept` and +// `iterator_category` +template +[[nodiscard]] _CCCL_API _CCCL_CONSTEVAL auto __iterator_traits_category_or_concept() noexcept +{ + if constexpr (!__has_member_iterator_concept<_Iter>) + { + return __iterator_category_type<_Iter>{}; + } + else if constexpr (__has_contiguous_traversal<_Iter>) + { + return contiguous_iterator_tag{}; + } + else if constexpr (__has_random_access_traversal<_Iter>) + { + return random_access_iterator_tag{}; + } + else if constexpr (__has_bidirectional_traversal<_Iter>) + { + return bidirectional_iterator_tag{}; + } + else if constexpr (__has_forward_traversal<_Iter>) + { + return forward_iterator_tag{}; + } + else if constexpr (__has_input_traversal<_Iter>) + { + return input_iterator_tag{}; + } + else // if constexpr (__has_member_iterator_category<_Iter>) + { + return typename _Iter::iterator_category{}; + } +} + +template +using __iterator_traits_category_or_concept_t = decltype(::cuda::std::__iterator_traits_category_or_concept<_Iter>()); + _CCCL_END_NAMESPACE_CUDA_STD #include diff --git a/libcudacxx/include/cuda/std/__iterator/move_iterator.h b/libcudacxx/include/cuda/std/__iterator/move_iterator.h index aefea4155af..a0221c1d5e2 100644 --- a/libcudacxx/include/cuda/std/__iterator/move_iterator.h +++ b/libcudacxx/include/cuda/std/__iterator/move_iterator.h @@ -63,9 +63,9 @@ template struct __move_iter_category_base<_Iter> { using iterator_category = - _If::iterator_category, random_access_iterator_tag>, + _If, random_access_iterator_tag>, random_access_iterator_tag, - typename iterator_traits<_Iter>::iterator_category>; + __iterator_traits_category_or_concept_t<_Iter>>; }; template @@ -85,9 +85,9 @@ template struct __move_iter_category_base<_Iter, enable_if_t<__has_iter_category>>> { using iterator_category = - _If::iterator_category, random_access_iterator_tag>, + _If, random_access_iterator_tag>, random_access_iterator_tag, - typename iterator_traits<_Iter>::iterator_category>; + __iterator_traits_category_or_concept_t<_Iter>>; }; template diff --git a/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h b/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h index fa265165dd0..1bf8562abd6 100644 --- a/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h +++ b/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h @@ -99,9 +99,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT reverse_iterator using iterator_type = _Iter; using iterator_category = - _If<__has_random_access_traversal<_Iter>, - random_access_iterator_tag, - typename iterator_traits<_Iter>::iterator_category>; + _If<__has_random_access_traversal<_Iter>, random_access_iterator_tag, __iterator_traits_category_or_concept_t<_Iter>>; using pointer = typename iterator_traits<_Iter>::pointer; using iterator_concept = _If, random_access_iterator_tag, bidirectional_iterator_tag>; using value_type = iter_value_t<_Iter>; diff --git a/libcudacxx/include/cuda/std/__simd/iterator.h b/libcudacxx/include/cuda/std/__simd/iterator.h index 7e4f4ca569a..3ae51402c7b 100644 --- a/libcudacxx/include/cuda/std/__simd/iterator.h +++ b/libcudacxx/include/cuda/std/__simd/iterator.h @@ -265,11 +265,11 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD template struct iterator_traits> { - using __iter = simd::__simd_iterator<_Vp>; - using iterator_concept = typename __iter::iterator_concept; - using iterator_category = typename __iter::iterator_category; - using value_type = typename __iter::value_type; - using difference_type = typename __iter::difference_type; + using _Iter = simd::__simd_iterator<_Vp>; + using iterator_concept = typename _Iter::iterator_concept; + using iterator_category = typename _Iter::iterator_category; + using value_type = typename _Iter::value_type; + using difference_type = typename _Iter::difference_type; using pointer = void; using reference = value_type; }; diff --git a/thrust/testing/catch2_test_iterator_traits.cu b/thrust/testing/catch2_test_iterator_traits.cu index 7e602213d32..8b7a3c4b4c6 100644 --- a/thrust/testing/catch2_test_iterator_traits.cu +++ b/thrust/testing/catch2_test_iterator_traits.cu @@ -1,6 +1,7 @@ #include #include #include +#include #include #include @@ -56,6 +57,15 @@ struct cuda_make_counting_iterator } }; +struct cuda_make_permutation_iterator +{ + template + auto operator()(Args&&... args) const + { + return cuda::make_permutation_iterator(cuda::std::forward(args)...); + } +}; + struct cuda_make_transform_iterator { template @@ -83,6 +93,15 @@ struct thrust_make_counting_iterator } }; +struct thrust_make_permutation_iterator +{ + template + auto operator()(Args&&... args) const + { + return thrust::make_permutation_iterator(cuda::std::forward(args)...); + } +}; + struct thrust_make_transform_iterator { template @@ -108,16 +127,18 @@ inline constexpr bool has_random_access_traversal = using cuda::std::__type_cartesian_product; using cuda::std::__type_list; -using make_counting_t = __type_list; -using make_transform_t = __type_list; -using make_zip_t = __type_list; -using make_it_t = __type_cartesian_product; +using make_counting_t = __type_list; +using make_transform_t = __type_list; +using make_zip_t = __type_list; +using make_permutation_t = __type_list; +using make_it_t = __type_cartesian_product; 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; + __type_at_c<3, TestType> make_permutation_iterator; auto counting_it = make_counting_iterator(0); STATIC_REQUIRE(is_same_v, thrust::any_system_tag>); @@ -130,6 +151,10 @@ TEMPLATE_LIST_TEST_CASE("iterator system and traversal propagation - any system" [[maybe_unused]] auto zip_it = make_zip_iterator(counting_it, transform_it); STATIC_REQUIRE(is_same_v, thrust::any_system_tag>); STATIC_REQUIRE(has_random_access_traversal); + + [[maybe_unused]] auto permutation_it = make_permutation_iterator(transform_it, transform_it); + STATIC_REQUIRE(is_same_v, thrust::any_system_tag>); + STATIC_REQUIRE(has_random_access_traversal); } auto expected_tag(thrust::device_vector) -> thrust::device_system_tag; diff --git a/thrust/thrust/iterator/iterator_traits.h b/thrust/thrust/iterator/iterator_traits.h index 34308c8ef85..58eb8191b4b 100644 --- a/thrust/thrust/iterator/iterator_traits.h +++ b/thrust/thrust/iterator/iterator_traits.h @@ -148,7 +148,7 @@ using iterator_difference_t //! Trait obtaining the iterator traversal category of an iterator type. template struct iterator_traversal - : detail::iterator_category_to_traversal::iterator_category> + : detail::iterator_category_to_traversal<::cuda::std::__iterator_traits_category_or_concept_t> {}; //! Alias to the iterator traversal category of an iterator type. @@ -165,7 +165,7 @@ struct iterator_system_impl template struct iterator_system_impl::iterator_category>> - : iterator_category_to_system::iterator_category> + : iterator_category_to_system<::cuda::std::__iterator_traits_category_or_concept_t> {}; } // namespace detail