Deprecate and replace thrust::cuda_cub iterators

thrust/thrust/system/cuda/detail/count.h

@@ -40,6 +40,7 @@
 #  include <thrust/system/cuda/config.h>
 
 #  include <thrust/distance.h>
+#  include <thrust/iterator/transform_iterator.h>
 #  include <thrust/system/cuda/detail/reduce.h>
 #  include <thrust/system/cuda/detail/util.h>
 
@@ -52,7 +53,7 @@ typename iterator_traits<InputIt>::difference_type _CCCL_HOST_DEVICE
 count_if(execution_policy<Derived>& policy, InputIt first, InputIt last, UnaryPred unary_pred)
 {
   using size_type       = typename iterator_traits<InputIt>::difference_type;
-  using flag_iterator_t = transform_input_iterator_t<size_type, InputIt, UnaryPred>;
+  using flag_iterator_t = transform_iterator<UnaryPred, InputIt, size_type, size_type>;
 
   return cuda_cub::reduce_n(
     policy, flag_iterator_t(first, unary_pred), thrust::distance(first, last), size_type(0), plus<size_type>());

thrust/thrust/system/cuda/detail/extrema.h

@@ -28,6 +28,8 @@
 
 #include <thrust/detail/config.h>
 
+#include <thrust/iterator/transform_iterator.h>
+
 #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
 #  pragma GCC system_header
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
@@ -370,10 +372,10 @@ element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, BinaryPr
 
   IndexType num_items = static_cast<IndexType>(thrust::distance(first, last));
 
-  using iterator_tuple = tuple<ItemsIt, counting_iterator_t<IndexType>>;
+  using iterator_tuple = tuple<ItemsIt, counting_iterator<IndexType>>;
   using zip_iterator   = zip_iterator<iterator_tuple>;
 
-  iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator_t<IndexType>(0));
+  iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator<IndexType>(0));
 
   using arg_min_t = ArgFunctor<InputType, IndexType, BinaryPred>;
   using T         = tuple<InputType, IndexType>;
@@ -443,15 +445,15 @@ minmax_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, B
 
      const auto num_items = static_cast<IndexType>(thrust::distance(first, last));
 
-     using iterator_tuple = tuple<ItemsIt, counting_iterator_t<IndexType>>;
+     using iterator_tuple = tuple<ItemsIt, counting_iterator<IndexType>>;
      using zip_iterator   = zip_iterator<iterator_tuple>;
 
-     iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator_t<IndexType>(0));
+     iterator_tuple iter_tuple = thrust::make_tuple(first, counting_iterator<IndexType>(0));
 
      using arg_minmax_t   = __extrema::arg_minmax_f<InputType, IndexType, BinaryPred>;
      using two_pairs_type = typename arg_minmax_t::two_pairs_type;
      using duplicate_t    = typename arg_minmax_t::duplicate_tuple;
-     using transform_t    = transform_input_iterator_t<two_pairs_type, zip_iterator, duplicate_t>;
+     using transform_t    = transform_iterator<duplicate_t, zip_iterator, two_pairs_type, two_pairs_type>;
 
      zip_iterator begin    = make_zip_iterator(iter_tuple);
      two_pairs_type result = __extrema::extrema(

thrust/thrust/system/cuda/detail/find.h

@@ -28,6 +28,8 @@
 
 #include <thrust/detail/config.h>
 
+#include <thrust/iterator/transform_iterator.h>
+
 #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
 #  pragma GCC system_header
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
@@ -115,11 +117,11 @@ find_if_n(execution_policy<Derived>& policy, InputIt first, Size num_items, Pred
   const Size interval_size      = (::cuda::std::min)(interval_threshold, num_items);
 
   // force transform_iterator output to bool
-  using XfrmIterator  = transform_input_iterator_t<bool, InputIt, Predicate>;
-  using IteratorTuple = thrust::tuple<XfrmIterator, counting_iterator_t<Size>>;
+  using XfrmIterator  = transform_iterator<Predicate, InputIt, bool, bool>;
+  using IteratorTuple = thrust::tuple<XfrmIterator, counting_iterator<Size>>;
   using ZipIterator   = thrust::zip_iterator<IteratorTuple>;
 
-  IteratorTuple iter_tuple = thrust::make_tuple(XfrmIterator(first, predicate), counting_iterator_t<Size>(0));
+  IteratorTuple iter_tuple = thrust::make_tuple(XfrmIterator(first, predicate), counting_iterator<Size>(0));
 
   ZipIterator begin = thrust::make_zip_iterator(iter_tuple);
   ZipIterator end   = begin + num_items;

thrust/thrust/system/cuda/detail/inner_product.h

@@ -38,15 +38,15 @@
 
 #if _CCCL_HAS_CUDA_COMPILER
 #  include <thrust/distance.h>
+#  include <thrust/iterator/transform_iterator.h>
+#  include <thrust/iterator/zip_iterator.h>
 #  include <thrust/system/cuda/detail/reduce.h>
-
-#  include <iterator>
+#  include <thrust/zip_function.h>
 
 THRUST_NAMESPACE_BEGIN
 
 namespace cuda_cub
 {
-
 template <class Derived, class InputIt1, class InputIt2, class T, class ReduceOp, class ProductOp>
 T _CCCL_HOST_DEVICE inner_product(
   execution_policy<Derived>& policy,
@@ -57,11 +57,9 @@ T _CCCL_HOST_DEVICE inner_product(
   ReduceOp reduce_op,
   ProductOp product_op)
 {
-  using size_type        = typename iterator_traits<InputIt1>::difference_type;
-  size_type num_items    = static_cast<size_type>(thrust::distance(first1, last1));
-  using binop_iterator_t = transform_pair_of_input_iterators_t<T, InputIt1, InputIt2, ProductOp>;
-
-  return cuda_cub::reduce_n(policy, binop_iterator_t(first1, first2, product_op), num_items, init, reduce_op);
+  const auto n     = thrust::distance(first1, last1);
+  const auto first = make_transform_iterator(make_zip_iterator(first1, first2), make_zip_function(product_op));
+  return cuda_cub::reduce_n(policy, first, n, init, reduce_op);
 }
 
 template <class Derived, class InputIt1, class InputIt2, class T>
@@ -70,7 +68,6 @@ inner_product(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1
 {
   return cuda_cub::inner_product(policy, first1, last1, first2, init, plus<T>(), multiplies<T>());
 }
-
 } // namespace cuda_cub
 
 THRUST_NAMESPACE_END

thrust/thrust/system/cuda/detail/mismatch.h

@@ -40,15 +40,14 @@
 #  include <thrust/system/cuda/config.h>
 
 #  include <thrust/distance.h>
+#  include <thrust/iterator/zip_iterator.h>
 #  include <thrust/pair.h>
 #  include <thrust/system/cuda/detail/execution_policy.h>
-
-#  include <cuda/std/__functional/identity.h>
+#  include <thrust/zip_function.h>
 
 THRUST_NAMESPACE_BEGIN
 namespace cuda_cub
 {
-
 template <class Derived, class InputIt1, class InputIt2, class BinaryPred>
 pair<InputIt1, InputIt2> _CCCL_HOST_DEVICE
 mismatch(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1, InputIt2 first2, BinaryPred binary_pred);
@@ -69,15 +68,12 @@ template <class Derived, class InputIt1, class InputIt2, class BinaryPred>
 pair<InputIt1, InputIt2> _CCCL_HOST_DEVICE
 mismatch(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1, InputIt2 first2, BinaryPred binary_pred)
 {
-  using transform_t = transform_pair_of_input_iterators_t<bool, InputIt1, InputIt2, BinaryPred>;
-
-  transform_t transform_first = transform_t(first1, first2, binary_pred);
-
-  transform_t result = cuda_cub::find_if_not(
-    policy, transform_first, transform_first + thrust::distance(first1, last1), ::cuda::std::__identity{});
-
-  return thrust::make_pair(first1 + thrust::distance(transform_first, result),
-                           first2 + thrust::distance(transform_first, result));
+  const auto n            = thrust::distance(first1, last1);
+  const auto first        = make_zip_iterator(first1, first2);
+  const auto last         = make_zip_iterator(last1, first2 + n);
+  const auto mismatch_pos = cuda_cub::find_if_not(policy, first, last, make_zip_function(binary_pred));
+  const auto dist         = thrust::distance(first, mismatch_pos);
+  return thrust::make_pair(first1 + dist, first2 + dist);
 }
 
 template <class Derived, class InputIt1, class InputIt2>

thrust/thrust/system/cuda/detail/transform_scan.h

@@ -28,6 +28,8 @@
 
 #include <thrust/detail/config.h>
 
+#include "thrust/iterator/transform_iterator.h"
+
 #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
 #  pragma GCC system_header
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
@@ -66,7 +68,7 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan(
 
   using size_type              = typename iterator_traits<InputIt>::difference_type;
   size_type num_items          = static_cast<size_type>(thrust::distance(first, last));
-  using transformed_iterator_t = transform_input_iterator_t<value_type, InputIt, TransformOp>;
+  using transformed_iterator_t = transform_iterator<TransformOp, InputIt, value_type, value_type>;
 
   return cuda_cub::inclusive_scan_n(policy, transformed_iterator_t(first, transform_op), num_items, result, scan_op);
 }
@@ -87,7 +89,7 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan(
 
   using size_type              = typename iterator_traits<InputIt>::difference_type;
   size_type num_items          = static_cast<size_type>(thrust::distance(first, last));
-  using transformed_iterator_t = transform_input_iterator_t<value_type, InputIt, TransformOp>;
+  using transformed_iterator_t = transform_iterator<TransformOp, InputIt, value_type, value_type>;
 
   return cuda_cub::inclusive_scan_n(
     policy, transformed_iterator_t(first, transform_op), num_items, result, init, scan_op);
@@ -108,7 +110,7 @@ OutputIt _CCCL_HOST_DEVICE transform_exclusive_scan(
 
   using size_type              = typename iterator_traits<InputIt>::difference_type;
   size_type num_items          = static_cast<size_type>(thrust::distance(first, last));
-  using transformed_iterator_t = transform_input_iterator_t<result_type, InputIt, TransformOp>;
+  using transformed_iterator_t = transform_iterator<TransformOp, InputIt, result_type, result_type>;
 
   return cuda_cub::exclusive_scan_n(
     policy, transformed_iterator_t(first, transform_op), num_items, result, init, scan_op);

thrust/thrust/system/cuda/detail/util.h

@@ -249,12 +249,13 @@ _CCCL_HOST_DEVICE inline void throw_on_error(cudaError_t status, char const* msg
   }
 }
 
-// FIXME: Move the iterators elsewhere.
-
+// deprecated [Since 2.8]
 template <class ValueType, class InputIt, class UnaryOp>
-struct transform_input_iterator_t
+struct CCCL_DEPRECATED_BECAUSE("Use thrust::transform_iterator") transform_input_iterator_t
 {
-  using self_t            = transform_input_iterator_t;
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
+  using self_t = transform_input_iterator_t;
+  _CCCL_SUPPRESS_DEPRECATED_POP
   using difference_type   = typename iterator_traits<InputIt>::difference_type;
   using value_type        = ValueType;
   using pointer           = void;
@@ -358,10 +359,14 @@ struct transform_input_iterator_t
   }
 }; // struct transform_input_iterarot_t
 
+// deprecated [Since 2.8]
 template <class ValueType, class InputIt1, class InputIt2, class BinaryOp>
-struct transform_pair_of_input_iterators_t
+struct CCCL_DEPRECATED_BECAUSE("Use thrust::transform_iterator of a thrust::zip_iterator")
+  transform_pair_of_input_iterators_t
 {
-  using self_t            = transform_pair_of_input_iterators_t;
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
+  using self_t = transform_pair_of_input_iterators_t;
+  _CCCL_SUPPRESS_DEPRECATED_POP
   using difference_type   = typename iterator_traits<InputIt1>::difference_type;
   using value_type        = ValueType;
   using pointer           = void;
@@ -488,10 +493,13 @@ struct CCCL_DEPRECATED_BECAUSE("Use cuda::std::identity") identity
   }
 };
 
+// deprecated [Since 2.8]
 template <class T>
-struct counting_iterator_t
+struct CCCL_DEPRECATED_BECAUSE("Use thrust::counting_iterator") counting_iterator_t
 {
-  using self_t            = counting_iterator_t;
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
+  using self_t = counting_iterator_t;
+  _CCCL_SUPPRESS_DEPRECATED_POP
   using difference_type   = T;
   using value_type        = T;
   using pointer           = void;