Replace thrust::discard_iterator with cuda::discard_iterator

cub/cub/device/device_transform.cuh

@@ -23,6 +23,7 @@
 #include <cuda/__iterator/zip_iterator.h>
 #include <cuda/__stream/get_stream.h>
 #include <cuda/std/__execution/env.h>
+#include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/tuple>
 
 CUB_NAMESPACE_BEGIN
@@ -433,7 +434,7 @@ struct DeviceTransform
   {
     static_assert(::cuda::std::is_invocable_v<Generator>, "The passed generator must be a nullary function object");
     static_assert(
-      ::cuda::std::is_assignable_v<detail::it_reference_t<RandomAccessIteratorOut>,
+      ::cuda::std::is_assignable_v<::cuda::std::iter_reference_t<RandomAccessIteratorOut>,
                                    ::cuda::std::invoke_result_t<Generator>>,
       "The return value of the generator's call operator must be assignable to the dereferenced output iterator");
 
@@ -499,7 +500,7 @@ struct DeviceTransform
   CUB_RUNTIME_FUNCTION static cudaError_t
   Fill(RandomAccessIteratorOut output, NumItemsT num_items, Value value, Env env = {})
   {
-    static_assert(::cuda::std::is_assignable_v<detail::it_reference_t<RandomAccessIteratorOut>, Value>,
+    static_assert(::cuda::std::is_assignable_v<::cuda::std::iter_reference_t<RandomAccessIteratorOut>, Value>,
                   "The passed value must be assignable to the dereferenced output iterator");
 
     _CCCL_NVTX_RANGE_SCOPE("cub::DeviceTransform::Fill");

cub/cub/device/dispatch/dispatch_scan.cuh

@@ -1029,20 +1029,21 @@ template <typename AccumT,
           typename ScanOpT,
           typename InitValueT,
           typename OffsetT,
-          typename PolicySelector = policy_selector_from_types<detail::it_value_t<InputIteratorT>,
-                                                               detail::it_value_t<OutputIteratorT>,
-                                                               AccumT,
-                                                               OffsetT,
-                                                               ScanOpT>,
-          typename KernelSource   = DeviceScanKernelSource<
-              PolicySelector,
-              THRUST_NS_QUALIFIER::try_unwrap_contiguous_iterator_t<InputIteratorT>,
-              THRUST_NS_QUALIFIER::try_unwrap_contiguous_iterator_t<OutputIteratorT>,
-              ScanOpT,
-              InitValueT,
-              OffsetT,
-              AccumT,
-              EnforceInclusive>,
+          typename PolicySelector =
+            policy_selector_from_types<detail::non_void_value_t<InputIteratorT, detail::it_value_t<OutputIteratorT>>,
+                                       detail::non_void_value_t<OutputIteratorT, detail::it_value_t<InputIteratorT>>,
+                                       AccumT,
+                                       OffsetT,
+                                       ScanOpT>,
+          typename KernelSource = DeviceScanKernelSource<
+            PolicySelector,
+            THRUST_NS_QUALIFIER::try_unwrap_contiguous_iterator_t<InputIteratorT>,
+            THRUST_NS_QUALIFIER::try_unwrap_contiguous_iterator_t<OutputIteratorT>,
+            ScanOpT,
+            InitValueT,
+            OffsetT,
+            AccumT,
+            EnforceInclusive>,
           typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY>
 CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch_with_accum(
   void* d_temp_storage,

cub/cub/device/dispatch/kernels/kernel_scan_warpspeed.cuh

@@ -24,6 +24,7 @@
 #include <cub/device/dispatch/kernels/scan_warpspeed_policy.cuh>
 #include <cub/thread/thread_reduce.cuh>
 #include <cub/thread/thread_scan.cuh>
+#include <cub/util_type.cuh>
 #include <cub/warp/warp_reduce.cuh>
 #include <cub/warp/warp_scan.cuh>
 
@@ -935,8 +936,8 @@ _CCCL_API constexpr bool use_warpspeed(
 template <typename InputIteratorT, typename OutputIteratorT, typename AccumT>
 _CCCL_API constexpr bool use_warpspeed(const scan_warpspeed_policy& policy)
 {
-  using InputT  = it_value_t<InputIteratorT>;
-  using OutputT = it_value_t<OutputIteratorT>;
+  using InputT  = non_void_value_t<InputIteratorT, it_value_t<OutputIteratorT>>;
+  using OutputT = non_void_value_t<OutputIteratorT, it_value_t<InputIteratorT>>;
   return use_warpspeed(
     policy,
     static_cast<int>(sizeof(InputT)),

cub/test/catch2_test_device_select_flagged_if.cu

@@ -5,11 +5,11 @@
 
 #include <cub/device/device_select.cuh>
 
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/logical.h>
 
 #include <cuda/functional>
+#include <cuda/iterator>
 
 #include <algorithm>
 
@@ -38,7 +38,7 @@ get_reference(c2h::device_vector<T> const& in, c2h::device_vector<FlagT> const&
   auto zipped_in_it = thrust::make_zip_iterator(h_flags.cbegin(), reference.cbegin());
 
   // Discards the flags part and only keeps the items
-  auto zipped_out_it = thrust::make_zip_iterator(thrust::make_discard_iterator(), reference.begin());
+  auto zipped_out_it = thrust::make_zip_iterator(cuda::make_discard_iterator(), reference.begin());
 
   auto end =
     std::copy_if(zipped_in_it, zipped_in_it + in.size(), zipped_out_it, predicate_op_wrapper_t<Pred>{if_predicate});

cub/test/catch2_test_nvrtc.cu

@@ -92,7 +92,7 @@ TEST_CASE("Test nvrtc", "[test][nvrtc]")
 
     #include <thrust/iterator/constant_iterator.h>
     #include <thrust/iterator/counting_iterator.h>
-    #include <thrust/iterator/discard_iterator.h>
+    #include <cuda/iterator>
     #include <thrust/iterator/permutation_iterator.h>
     #include <thrust/iterator/reverse_iterator.h>
     #include <thrust/iterator/tabulate_output_iterator.h>

cub/test/test_util.h

@@ -19,8 +19,7 @@
 #include <cub/util_ptx.cuh>
 #include <cub/util_type.cuh>
 
-#include <thrust/iterator/discard_iterator.h>
-
+#include <cuda/iterator>
 #include <cuda/std/algorithm>
 
 #include <nv/target>
@@ -1150,10 +1149,10 @@ inline int CompareDeviceResults(
  * Verify the contents of a device array match those
  * of a host array
  */
-template <typename S, typename OffsetT>
+template <typename S>
 int CompareDeviceResults(
   S* /*h_reference*/,
-  THRUST_NS_QUALIFIER::discard_iterator<OffsetT> /*d_data*/,
+  cuda::discard_iterator /*d_data*/,
   std::size_t /*num_items*/,
   bool /*verbose*/      = true,
   bool /*display_data*/ = false)

libcudacxx/include/cuda/__iterator/discard_iterator.h

@@ -25,6 +25,7 @@
 #include <cuda/std/__iterator/concepts.h>
 #include <cuda/std/__iterator/default_sentinel.h>
 #include <cuda/std/__iterator/iterator_traits.h>
+#include <cuda/std/__iterator/readable_traits.h>
 #include <cuda/std/__type_traits/is_same.h>
 #include <cuda/std/__type_traits/remove_cvref.h>
 #include <cuda/std/cstdint>
@@ -76,6 +77,12 @@ class discard_iterator
 public:
   struct __discard_proxy
   {
+    constexpr __discard_proxy() noexcept = default;
+
+    template <class _Tp>
+    _CCCL_API constexpr explicit __discard_proxy(const _Tp&) noexcept
+    {}
+
     _CCCL_TEMPLATE(class _Tp)
     _CCCL_REQUIRES((!::cuda::std::is_same_v<::cuda::std::remove_cvref_t<_Tp>, __discard_proxy>) )
     _CCCL_API constexpr const __discard_proxy& operator=(_Tp&&) const noexcept
@@ -320,6 +327,20 @@ _CCCL_REQUIRES(::cuda::std::__integer_like<_Integer>)
 
 _CCCL_END_NAMESPACE_CUDA
 
+_CCCL_BEGIN_NAMESPACE_CUDA_STD
+
+// We need to specialize these traits because discard_iterator's value_type is usually
+// void, but you are allowed to read (and write) to it nonetheless. If we don't specialize
+// here, then quite a few downstream algorithms (like zip's) break because they only see a
+// value_type of void.
+template <>
+struct indirectly_readable_traits<::cuda::discard_iterator>
+{
+  using value_type = ::cuda::discard_iterator::__discard_proxy;
+};
+
+_CCCL_END_NAMESPACE_CUDA_STD
+
 #include <cuda/std/__cccl/epilogue.h>
 
 #endif // _CUDA___ITERATOR_DISCARD_ITERATOR_H

thrust/cmake/ThrustHeaderTesting.cmake

@@ -10,6 +10,7 @@ set(
   "thrust/iterator/tabulate_output_iterator\\.h"
   "thrust/iterator/strided_iterator\\.h"
   "thrust/iterator/constant_iterator\\.h"
+  "thrust/iterator/discard_iterator\\.h"
 )
 
 cccl_get_cudatoolkit()

thrust/examples/set_operations.cu

@@ -1,9 +1,10 @@
 #include <thrust/device_vector.h>
 #include <thrust/extrema.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/merge.h>
 #include <thrust/set_operations.h>
 
+#include <cuda/iterator>
+
 #include <iostream>
 
 // This example illustrates use of the set operation algorithms
@@ -127,7 +128,7 @@ template <typename Vector>
 void SetIntersectionSize(const Vector& A, const Vector& B)
 {
   // computes the exact size of the intersection without allocating output
-  thrust::discard_iterator<> C_begin, C_end;
+  cuda::discard_iterator C_begin, C_end;
 
   C_end = thrust::set_intersection(A.begin(), A.end(), B.begin(), B.end(), C_begin);
 

thrust/examples/sum_columns.cu

@@ -1,7 +1,6 @@
 #include <thrust/execution_policy.h>
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/random.h>
 #include <thrust/reduce.h>
@@ -53,7 +52,7 @@ int main()
   // Sum each column, storing the result in a new vector.
   thrust::universal_vector<int> sums(cols);
   thrust::reduce_by_key(
-    thrust::device, col_idx_begin, col_idx_end, M_transposed, thrust::make_discard_iterator(), sums.begin());
+    thrust::device, col_idx_begin, col_idx_end, M_transposed, cuda::make_discard_iterator(), sums.begin());
 
   // Output the result.
   thrust::for_each_n(thrust::seq, flat_idx, rows, [&](int i) {

thrust/examples/sum_rows.cu

@@ -1,7 +1,6 @@
 #include <thrust/execution_policy.h>
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/random.h>
 #include <thrust/reduce.h>
@@ -45,7 +44,7 @@ int main()
   // Sum each row, storing the result in a new vector.
   thrust::universal_vector<int> sums(rows);
   thrust::reduce_by_key(
-    thrust::device, row_idx_begin, row_idx_end, M.data_handle(), thrust::make_discard_iterator(), sums.begin());
+    thrust::device, row_idx_begin, row_idx_end, M.data_handle(), cuda::make_discard_iterator(), sums.begin());
 
   // Output the result.
   thrust::for_each_n(thrust::seq, flat_idx, rows, [&](int i) {

thrust/testing/binary_search_vector.cu

@@ -1,10 +1,11 @@
 #include <thrust/binary_search.h>
 #include <thrust/detail/allocator/allocator_system.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/sequence.h>
 #include <thrust/sort.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 //////////////////////
@@ -343,12 +344,12 @@ struct TestVectorLowerBoundDiscardIterator
     thrust::host_vector<T> h_input   = unittest::random_integers<T>(2 * n);
     thrust::device_vector<T> d_input = h_input;
 
-    thrust::discard_iterator<> h_result =
-      thrust::lower_bound(h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), thrust::make_discard_iterator());
-    thrust::discard_iterator<> d_result =
-      thrust::lower_bound(d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator h_result =
+      thrust::lower_bound(h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), cuda::make_discard_iterator());
+    cuda::discard_iterator d_result =
+      thrust::lower_bound(d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> reference(2 * n);
+    cuda::discard_iterator reference(2 * n);
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
@@ -368,12 +369,12 @@ struct TestVectorUpperBoundDiscardIterator
     thrust::host_vector<T> h_input   = unittest::random_integers<T>(2 * n);
     thrust::device_vector<T> d_input = h_input;
 
-    thrust::discard_iterator<> h_result =
-      thrust::upper_bound(h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), thrust::make_discard_iterator());
-    thrust::discard_iterator<> d_result =
-      thrust::upper_bound(d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator h_result =
+      thrust::upper_bound(h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), cuda::make_discard_iterator());
+    cuda::discard_iterator d_result =
+      thrust::upper_bound(d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> reference(2 * n);
+    cuda::discard_iterator reference(2 * n);
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
@@ -393,12 +394,12 @@ struct TestVectorBinarySearchDiscardIterator
     thrust::host_vector<T> h_input   = unittest::random_integers<T>(2 * n);
     thrust::device_vector<T> d_input = h_input;
 
-    thrust::discard_iterator<> h_result = thrust::binary_search(
-      h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), thrust::make_discard_iterator());
-    thrust::discard_iterator<> d_result = thrust::binary_search(
-      d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator h_result =
+      thrust::binary_search(h_vec.begin(), h_vec.end(), h_input.begin(), h_input.end(), cuda::make_discard_iterator());
+    cuda::discard_iterator d_result =
+      thrust::binary_search(d_vec.begin(), d_vec.end(), d_input.begin(), d_input.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> reference(2 * n);
+    cuda::discard_iterator reference(2 * n);
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);

thrust/testing/catch2_test_adjacent_difference.cu

@@ -1,9 +1,10 @@
 #include <thrust/adjacent_difference.h>
 #include <thrust/device_free.h>
 #include <thrust/device_malloc.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 
+#include <cuda/iterator>
+
 #include "catch2_test_helper.h"
 #include <unittest/random.h>
 #include <unittest/special_types.h>
@@ -112,12 +113,12 @@ TEMPLATE_LIST_TEST_CASE("AdjacentDifferenceDiscardIterator", "[adjacent_differen
     thrust::host_vector<T> h_input   = unittest::random_samples<T>(n);
     thrust::device_vector<T> d_input = h_input;
 
-    thrust::discard_iterator<> h_result =
-      thrust::adjacent_difference(h_input.begin(), h_input.end(), thrust::make_discard_iterator());
-    thrust::discard_iterator<> d_result =
-      thrust::adjacent_difference(d_input.begin(), d_input.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator h_result =
+      thrust::adjacent_difference(h_input.begin(), h_input.end(), cuda::make_discard_iterator());
+    cuda::discard_iterator d_result =
+      thrust::adjacent_difference(d_input.begin(), d_input.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
     CHECK((reference == h_result));
     CHECK((reference == d_result));