[libcu++] Rework our handling of iterator_category

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 <class _InputIterator1, class _InputIterator2>
@@ -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

libcudacxx/include/cuda/std/__algorithm/fill.h

@@ -48,7 +48,7 @@ __fill(_RandomAccessIterator __first, _RandomAccessIterator __last, const _Tp& _
 template <class _ForwardIterator, class _Tp>
 _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

libcudacxx/include/cuda/std/__algorithm/find_end.h

@@ -207,8 +207,8 @@ template <class _ForwardIterator1, class _ForwardIterator2, class _BinaryPredica
     __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>{});
 }
 
 template <class _ForwardIterator1, class _ForwardIterator2>

libcudacxx/include/cuda/std/__algorithm/iterator_operations.h

@@ -55,9 +55,6 @@ struct _IterOps<_RangeAlgPolicy>
   template <class _Iter>
   using __value_type = iter_value_t<_Iter>;
 
-  template <class _Iter>
-  using __iterator_category = ::cuda::std::ranges::__iterator_concept<_Iter>;
-
   template <class _Iter>
   using __difference_type = iter_difference_t<_Iter>;
 
@@ -79,9 +76,6 @@ struct _IterOps<_ClassicAlgPolicy>
   template <class _Iter>
   using __value_type = typename iterator_traits<_Iter>::value_type;
 
-  template <class _Iter>
-  using __iterator_category = typename iterator_traits<_Iter>::iterator_category;
-
   template <class _Iter>
   using __difference_type = typename iterator_traits<_Iter>::difference_type;
 

libcudacxx/include/cuda/std/__algorithm/partition.h

@@ -107,7 +107,7 @@ _CCCL_EXEC_CHECK_DISABLE
 template <class _ForwardIterator, class _Predicate>
 _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;

libcudacxx/include/cuda/std/__algorithm/reverse.h

@@ -63,7 +63,7 @@ _CCCL_EXEC_CHECK_DISABLE
 template <class _AlgPolicy, class _BidirectionalIterator, class _Sentinel>
 _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());
 }
 

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());
 

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");

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;
 }
 

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 <class _ForwardIterator, class _Size, class _Tp>

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());
 }
 

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());
 }
 

libcudacxx/include/cuda/std/__algorithm/stable_partition.h

@@ -347,7 +347,7 @@ _CCCL_API _ForwardIterator __stable_partition(
 template <class _ForwardIterator, class _Predicate>
 _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());
 }

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:

libcudacxx/include/cuda/std/__iterator/iterator_traits.h

@@ -238,6 +238,9 @@ _CCCL_CONCEPT __has_member_pointer = _CCCL_REQUIRES_EXPR((_Tp))(typename(typenam
 template <class _Tp>
 _CCCL_CONCEPT __has_member_iterator_category = _CCCL_REQUIRES_EXPR((_Tp))(typename(typename _Tp::iterator_category));
 
+template <class _Tp>
+_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 <class _Iter>
 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 <class _Iter>
+[[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 <class _Iter>
+using __iterator_traits_category_or_concept_t = decltype(::cuda::std::__iterator_traits_category_or_concept<_Iter>());
+
 _CCCL_END_NAMESPACE_CUDA_STD
 
 #include <cuda/std/__cccl/epilogue.h>

libcudacxx/include/cuda/std/__iterator/move_iterator.h

@@ -63,9 +63,9 @@ template <class _Iter>
 struct __move_iter_category_base<_Iter>
 {
   using iterator_category =
-    _If<derived_from<typename iterator_traits<_Iter>::iterator_category, random_access_iterator_tag>,
+    _If<derived_from<__iterator_traits_category_or_concept_t<_Iter>, random_access_iterator_tag>,
         random_access_iterator_tag,
-        typename iterator_traits<_Iter>::iterator_category>;
+        __iterator_traits_category_or_concept_t<_Iter>>;
 };
 
 template <class _Iter, class _Sent>
@@ -85,9 +85,9 @@ template <class _Iter>
 struct __move_iter_category_base<_Iter, enable_if_t<__has_iter_category<iterator_traits<_Iter>>>>
 {
   using iterator_category =
-    _If<derived_from<typename iterator_traits<_Iter>::iterator_category, random_access_iterator_tag>,
+    _If<derived_from<__iterator_traits_category_or_concept_t<_Iter>, random_access_iterator_tag>,
         random_access_iterator_tag,
-        typename iterator_traits<_Iter>::iterator_category>;
+        __iterator_traits_category_or_concept_t<_Iter>>;
 };
 
 template <class _Iter, class _Sent>

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<_Iter>, random_access_iterator_tag, bidirectional_iterator_tag>;
   using value_type       = iter_value_t<_Iter>;

libcudacxx/include/cuda/std/__simd/iterator.h

@@ -265,11 +265,11 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD
 template <typename _Vp>
 struct iterator_traits<simd::__simd_iterator<_Vp>>
 {
-  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;
 };

thrust/testing/catch2_test_iterator_traits.cu

@@ -1,6 +1,7 @@
 #include <thrust/device_vector.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/iterator_traits.h>
+#include <thrust/iterator/permutation_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 
@@ -56,6 +57,15 @@ struct cuda_make_counting_iterator
   }
 };
 
+struct cuda_make_permutation_iterator
+{
+  template <typename... Args>
+  auto operator()(Args&&... args) const
+  {
+    return cuda::make_permutation_iterator(cuda::std::forward<Args>(args)...);
+  }
+};
+
 struct cuda_make_transform_iterator
 {
   template <typename... Args>
@@ -83,6 +93,15 @@ struct thrust_make_counting_iterator
   }
 };
 
+struct thrust_make_permutation_iterator
+{
+  template <typename... Args>
+  auto operator()(Args&&... args) const
+  {
+    return thrust::make_permutation_iterator(cuda::std::forward<Args>(args)...);
+  }
+};
+
 struct thrust_make_transform_iterator
 {
   template <typename... Args>
@@ -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<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>;
+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_permutation_t = __type_list<cuda_make_permutation_iterator, thrust_make_permutation_iterator>;
+using make_it_t          = __type_cartesian_product<make_counting_t, make_transform_t, make_zip_t, make_permutation_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;
+  __type_at_c<3, TestType> make_permutation_iterator;
 
   auto counting_it = make_counting_iterator(0);
   STATIC_REQUIRE(is_same_v<thrust::iterator_system_t<decltype(counting_it)>, 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::iterator_system_t<decltype(zip_it)>, thrust::any_system_tag>);
   STATIC_REQUIRE(has_random_access_traversal<decltype(zip_it)>);
+
+  [[maybe_unused]] auto permutation_it = make_permutation_iterator(transform_it, transform_it);
+  STATIC_REQUIRE(is_same_v<thrust::iterator_system_t<decltype(permutation_it)>, thrust::any_system_tag>);
+  STATIC_REQUIRE(has_random_access_traversal<decltype(permutation_it)>);
 }
 
 auto expected_tag(thrust::device_vector<int>) -> thrust::device_system_tag;

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 <typename Iterator>
 struct iterator_traversal
-    : detail::iterator_category_to_traversal<typename iterator_traits<Iterator>::iterator_category>
+    : detail::iterator_category_to_traversal<::cuda::std::__iterator_traits_category_or_concept_t<Iterator>>
 {};
 
 //! Alias to the iterator traversal category of an iterator type.
@@ -165,7 +165,7 @@ struct iterator_system_impl
 
 template <typename Iterator>
 struct iterator_system_impl<Iterator, ::cuda::std::void_t<typename iterator_traits<Iterator>::iterator_category>>
-    : iterator_category_to_system<typename iterator_traits<Iterator>::iterator_category>
+    : iterator_category_to_system<::cuda::std::__iterator_traits_category_or_concept_t<Iterator>>
 {};
 } // namespace detail