From 0d68108500dee7df7cceed479a321daabf9c2290 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Thu, 9 Jan 2025 10:56:46 +0100
Subject: [PATCH] Deprecate CUB iterators existing in Thrust

Fixes: #3261
---
 cub/cub/iterator/constant_input_iterator.cuh  |  2 +-
 cub/cub/iterator/counting_input_iterator.cuh  |  2 +-
 cub/cub/iterator/discard_output_iterator.cuh  |  2 +-
 cub/cub/iterator/transform_input_iterator.cuh |  2 +-
 cub/test/catch2_test_iterator.cu              | 12 +++++++++++-
 5 files changed, 15 insertions(+), 5 deletions(-)

diff --git a/cub/cub/iterator/constant_input_iterator.cuh b/cub/cub/iterator/constant_input_iterator.cuh
index 87252292d86..5f8a46f0e0b 100644
--- a/cub/cub/iterator/constant_input_iterator.cuh
+++ b/cub/cub/iterator/constant_input_iterator.cuh
@@ -87,7 +87,7 @@ CUB_NAMESPACE_BEGIN
  *   The difference type of this iterator (Default: @p ptrdiff_t)
  */
 template <typename ValueType, typename OffsetT = ptrdiff_t>
-class ConstantInputIterator
+class CCCL_DEPRECATED_BECAUSE("Use thrust::constant_iterator instead") ConstantInputIterator
 {
 public:
   // Required iterator traits
diff --git a/cub/cub/iterator/counting_input_iterator.cuh b/cub/cub/iterator/counting_input_iterator.cuh
index 529d6a990b1..ec185fa318e 100644
--- a/cub/cub/iterator/counting_input_iterator.cuh
+++ b/cub/cub/iterator/counting_input_iterator.cuh
@@ -90,7 +90,7 @@ CUB_NAMESPACE_BEGIN
  *   The difference type of this iterator (Default: @p ptrdiff_t)
  */
 template <typename ValueType, typename OffsetT = ptrdiff_t>
-class CountingInputIterator
+class CCCL_DEPRECATED_BECAUSE("Use thrust::counting_iterator instead") CountingInputIterator
 {
 public:
   // Required iterator traits
diff --git a/cub/cub/iterator/discard_output_iterator.cuh b/cub/cub/iterator/discard_output_iterator.cuh
index 4b3698d53f3..e350733850d 100644
--- a/cub/cub/iterator/discard_output_iterator.cuh
+++ b/cub/cub/iterator/discard_output_iterator.cuh
@@ -54,7 +54,7 @@ CUB_NAMESPACE_BEGIN
  * @brief A discard iterator
  */
 template <typename OffsetT = ptrdiff_t>
-class DiscardOutputIterator
+class CCCL_DEPRECATED_BECAUSE("Use thrust::discard_iterator instead") DiscardOutputIterator
 {
 public:
   // Required iterator traits
diff --git a/cub/cub/iterator/transform_input_iterator.cuh b/cub/cub/iterator/transform_input_iterator.cuh
index 9e3166afb82..f3bac10df29 100644
--- a/cub/cub/iterator/transform_input_iterator.cuh
+++ b/cub/cub/iterator/transform_input_iterator.cuh
@@ -110,7 +110,7 @@ CUB_NAMESPACE_BEGIN
  *   The difference type of this iterator (Default: @p ptrdiff_t)
  */
 template <typename ValueType, typename ConversionOp, typename InputIteratorT, typename OffsetT = ptrdiff_t>
-class TransformInputIterator
+class CCCL_DEPRECATED_BECAUSE("Use thrust::transform_iterator instead") TransformInputIterator
 {
 public:
   // Required iterator traits
diff --git a/cub/test/catch2_test_iterator.cu b/cub/test/catch2_test_iterator.cu
index a5dd52dfe04..cacf6e9de02 100644
--- a/cub/test/catch2_test_iterator.cu
+++ b/cub/test/catch2_test_iterator.cu
@@ -91,8 +91,10 @@ __global__ void test_iterator_kernel(InputIteratorT d_in, T* d_out, InputIterato
   d_itrs[1] = d_in; // Iterator at offset 0
 }
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
 template <typename InputIteratorT, typename T>
-void test_iterator(InputIteratorT d_in, const c2h::host_vector<T>& h_reference)
+void test_iterator(InputIteratorT d_in, const c2h::host_vector<T>& h_reference) //
+  _CCCL_SUPPRESS_DEPRECATED_POP
 {
   c2h::device_vector<T> d_out(h_reference.size());
   c2h::device_vector<InputIteratorT> d_itrs(2, d_in); // TODO(bgruber): using a raw allocation halves the compile time
@@ -113,7 +115,9 @@ C2H_TEST("Test constant iterator", "[iterator]", scalar_types)
   using T                = c2h::get<0, TestType>;
   const T base           = static_cast<T>(GENERATE(0, 99));
   const auto h_reference = c2h::host_vector<T>{base, base, base, base, base, base, base, base};
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   test_iterator(cub::ConstantInputIterator<T>(base), h_reference);
+  _CCCL_SUPPRESS_DEPRECATED_POP
 }
 
 C2H_TEST("Test counting iterator", "[iterator]", scalar_types)
@@ -129,7 +133,9 @@ C2H_TEST("Test counting iterator", "[iterator]", scalar_types)
     static_cast<T>(base + 21),
     static_cast<T>(base + 11),
     static_cast<T>(base + 0)};
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   test_iterator(cub::CountingInputIterator<T>(base), h_reference);
+  _CCCL_SUPPRESS_DEPRECATED_POP
 }
 
 using cache_modifiers =
@@ -187,9 +193,11 @@ C2H_TEST("Test transform iterator", "[iterator]", types)
     op(h_data[21]),
     op(h_data[11]),
     op(h_data[0])};
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   test_iterator(cub::TransformInputIterator<T, transform_op_t<T>, const T*>(
                   const_cast<const T*>(const_cast<const T*>(thrust::raw_pointer_cast(d_data.data()))), op),
                 h_reference);
+  _CCCL_SUPPRESS_DEPRECATED_POP
 }
 
 C2H_TEST("Test tex-obj texture iterator", "[iterator]", types)
@@ -233,7 +241,9 @@ C2H_TEST("Test texture transform iterator", "[iterator]", types)
   TextureIterator d_tex_itr;
   CubDebugExit(
     d_tex_itr.BindTexture(const_cast<const T*>(thrust::raw_pointer_cast(d_data.data())), sizeof(T) * TEST_VALUES));
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   cub::TransformInputIterator<T, transform_op_t<T>, TextureIterator> xform_itr(d_tex_itr, op);
+  _CCCL_SUPPRESS_DEPRECATED_POP
   test_iterator(xform_itr, h_reference);
   CubDebugExit(d_tex_itr.UnbindTexture());
 }