diff --git a/cub/cub/device/device_transform.cuh b/cub/cub/device/device_transform.cuh
index 1b21230619d..08fa9153d37 100644
--- a/cub/cub/device/device_transform.cuh
+++ b/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");
diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh
index 9d2b1cec65d..d86c3b63ad0 100644
--- a/cub/cub/device/dispatch/dispatch_scan.cuh
+++ b/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,
diff --git a/cub/cub/device/dispatch/kernels/kernel_scan_warpspeed.cuh b/cub/cub/device/dispatch/kernels/kernel_scan_warpspeed.cuh
index 09c88012563..42c8f3372ab 100644
--- a/cub/cub/device/dispatch/kernels/kernel_scan_warpspeed.cuh
+++ b/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)),
diff --git a/cub/test/catch2_test_device_select_flagged_if.cu b/cub/test/catch2_test_device_select_flagged_if.cu
index 686d1467f4d..5c9065e0a29 100644
--- a/cub/test/catch2_test_device_select_flagged_if.cu
+++ b/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});
diff --git a/cub/test/catch2_test_nvrtc.cu b/cub/test/catch2_test_nvrtc.cu
index fbc9f7c668a..6b7018c9fcb 100644
--- a/cub/test/catch2_test_nvrtc.cu
+++ b/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>
diff --git a/cub/test/test_util.h b/cub/test/test_util.h
index 8ac38867f49..19ddb6bdc45 100644
--- a/cub/test/test_util.h
+++ b/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)
diff --git a/libcudacxx/include/cuda/__iterator/discard_iterator.h b/libcudacxx/include/cuda/__iterator/discard_iterator.h
index 15de4d23f7d..0013efd2a36 100644
--- a/libcudacxx/include/cuda/__iterator/discard_iterator.h
+++ b/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
diff --git a/thrust/cmake/ThrustHeaderTesting.cmake b/thrust/cmake/ThrustHeaderTesting.cmake
index a3503b308f5..72eba9956be 100644
--- a/thrust/cmake/ThrustHeaderTesting.cmake
+++ b/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()
diff --git a/thrust/examples/set_operations.cu b/thrust/examples/set_operations.cu
index 0942248cf5c..501aadb4536 100644
--- a/thrust/examples/set_operations.cu
+++ b/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);
 
diff --git a/thrust/examples/sum_columns.cu b/thrust/examples/sum_columns.cu
index 401990f4357..cbe6b8e801f 100644
--- a/thrust/examples/sum_columns.cu
+++ b/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) {
diff --git a/thrust/examples/sum_rows.cu b/thrust/examples/sum_rows.cu
index 84123b79321..0e956140d04 100644
--- a/thrust/examples/sum_rows.cu
+++ b/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) {
diff --git a/thrust/testing/binary_search_vector.cu b/thrust/testing/binary_search_vector.cu
index d9faf363b30..3a3b2c3287a 100644
--- a/thrust/testing/binary_search_vector.cu
+++ b/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);
diff --git a/thrust/testing/catch2_test_adjacent_difference.cu b/thrust/testing/catch2_test_adjacent_difference.cu
index 06eabf9bb74..ce412a48283 100644
--- a/thrust/testing/catch2_test_adjacent_difference.cu
+++ b/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));
diff --git a/thrust/testing/catch2_test_transform.cu b/thrust/testing/catch2_test_transform.cu
index ec45eddbdc4..2adb0a8c245 100644
--- a/thrust/testing/catch2_test_transform.cu
+++ b/thrust/testing/catch2_test_transform.cu
@@ -1,9 +1,9 @@
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform.h>
 
+#include <cuda/iterator>
 #include <cuda/std/tuple>
 #include <cuda/std/utility>
 
@@ -361,13 +361,13 @@ TEMPLATE_LIST_TEST_CASE("UnaryToDiscardIterator", "[transform]", variable_list)
     thrust::host_vector<T> h_input   = unittest::random_integers<T>(n);
     thrust::device_vector<T> d_input = h_input;
 
-    thrust::discard_iterator<> h_result =
-      thrust::transform(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), ::cuda::std::negate<T>());
+    cuda::discard_iterator h_result =
+      thrust::transform(h_input.begin(), h_input.end(), cuda::make_discard_iterator(), ::cuda::std::negate<T>());
 
-    thrust::discard_iterator<> d_result =
-      thrust::transform(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), ::cuda::std::negate<T>());
+    cuda::discard_iterator d_result =
+      thrust::transform(d_input.begin(), d_input.end(), cuda::make_discard_iterator(), ::cuda::std::negate<T>());
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
     CHECK((reference == h_result));
     CHECK((reference == d_result));
@@ -397,20 +397,20 @@ TEMPLATE_LIST_TEST_CASE("UnaryToDiscardIteratorZipped", "[transform]", variable_
     using Iterator1 = typename thrust::host_vector<T>::iterator;
     using Iterator2 = typename thrust::device_vector<T>::iterator;
 
-    using Tuple1 = cuda::std::tuple<Iterator1, thrust::discard_iterator<>>;
-    using Tuple2 = cuda::std::tuple<Iterator2, thrust::discard_iterator<>>;
+    using Tuple1 = cuda::std::tuple<Iterator1, cuda::discard_iterator>;
+    using Tuple2 = cuda::std::tuple<Iterator2, cuda::discard_iterator>;
 
     using ZipIterator1 = thrust::zip_iterator<Tuple1>;
     using ZipIterator2 = thrust::zip_iterator<Tuple2>;
 
-    ZipIterator1 z1(cuda::std::tuple(h_output.begin(), thrust::make_discard_iterator()));
-    ZipIterator2 z2(cuda::std::tuple(d_output.begin(), thrust::make_discard_iterator()));
+    ZipIterator1 z1(cuda::std::tuple(h_output.begin(), cuda::make_discard_iterator()));
+    ZipIterator2 z2(cuda::std::tuple(d_output.begin(), cuda::make_discard_iterator()));
 
     ZipIterator1 h_result = thrust::transform(h_input.begin(), h_input.end(), z1, repeat2());
 
     ZipIterator2 d_result = thrust::transform(d_input.begin(), d_input.end(), z2, repeat2());
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
     CHECK(h_output == d_output);
 
@@ -480,23 +480,23 @@ TEMPLATE_LIST_TEST_CASE("UnaryToDiscardIterator", "[transform_if]", variable_lis
     thrust::device_vector<T> d_input   = h_input;
     thrust::device_vector<T> d_stencil = h_stencil;
 
-    thrust::discard_iterator<> h_result = thrust::transform_if(
+    cuda::discard_iterator h_result = thrust::transform_if(
       h_input.begin(),
       h_input.end(),
       h_stencil.begin(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       ::cuda::std::negate<T>(),
       is_positive());
 
-    thrust::discard_iterator<> d_result = thrust::transform_if(
+    cuda::discard_iterator d_result = thrust::transform_if(
       d_input.begin(),
       d_input.end(),
       d_stencil.begin(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       ::cuda::std::negate<T>(),
       is_positive());
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
     CHECK((reference == h_result));
     CHECK((reference == d_result));
@@ -540,12 +540,12 @@ TEMPLATE_LIST_TEST_CASE("BinaryToDiscardIterator", "[transform]", variable_list)
     thrust::device_vector<T> d_input1 = h_input1;
     thrust::device_vector<T> d_input2 = h_input2;
 
-    thrust::discard_iterator<> h_result = thrust::transform(
-      h_input1.begin(), h_input1.end(), h_input2.begin(), thrust::make_discard_iterator(), ::cuda::std::minus<T>());
-    thrust::discard_iterator<> d_result = thrust::transform(
-      d_input1.begin(), d_input1.end(), d_input2.begin(), thrust::make_discard_iterator(), ::cuda::std::minus<T>());
+    cuda::discard_iterator h_result = thrust::transform(
+      h_input1.begin(), h_input1.end(), h_input2.begin(), cuda::make_discard_iterator(), ::cuda::std::minus<T>());
+    cuda::discard_iterator d_result = thrust::transform(
+      d_input1.begin(), d_input1.end(), d_input2.begin(), cuda::make_discard_iterator(), ::cuda::std::minus<T>());
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
     CHECK((reference == h_result));
     CHECK((reference == d_result));
@@ -625,25 +625,25 @@ TEMPLATE_LIST_TEST_CASE("BinaryToDiscardIterator", "[transform_if]", variable_li
     thrust::device_vector<T> d_input2  = h_input2;
     thrust::device_vector<T> d_stencil = h_stencil;
 
-    thrust::discard_iterator<> h_result = thrust::transform_if(
+    cuda::discard_iterator h_result = thrust::transform_if(
       h_input1.begin(),
       h_input1.end(),
       h_input2.begin(),
       h_stencil.begin(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       ::cuda::std::minus<T>(),
       is_positive());
 
-    thrust::discard_iterator<> d_result = thrust::transform_if(
+    cuda::discard_iterator d_result = thrust::transform_if(
       d_input1.begin(),
       d_input1.end(),
       d_input2.begin(),
       d_stencil.begin(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       ::cuda::std::minus<T>(),
       is_positive());
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
     CHECK((reference == h_result));
     CHECK((reference == d_result));
diff --git a/thrust/testing/copy.cu b/thrust/testing/copy.cu
index 73a6892d66e..16b714b690e 100644
--- a/thrust/testing/copy.cu
+++ b/thrust/testing/copy.cu
@@ -4,7 +4,6 @@
 #include <thrust/device_free.h>
 #include <thrust/device_malloc.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/sequence.h>
@@ -57,13 +56,13 @@ void TestCopyToDiscardIterator()
   thrust::host_vector<T> h_input(5, 1);
   thrust::device_vector<T> d_input = h_input;
 
-  thrust::discard_iterator<> reference(5);
+  cuda::discard_iterator reference(5);
 
   // copy from host_vector
-  thrust::discard_iterator<> h_result = thrust::copy(h_input.begin(), h_input.end(), thrust::make_discard_iterator());
+  cuda::discard_iterator h_result = thrust::copy(h_input.begin(), h_input.end(), cuda::make_discard_iterator());
 
   // copy from device_vector
-  thrust::discard_iterator<> d_result = thrust::copy(d_input.begin(), d_input.end(), thrust::make_discard_iterator());
+  cuda::discard_iterator d_result = thrust::copy(d_input.begin(), d_input.end(), cuda::make_discard_iterator());
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
@@ -79,10 +78,10 @@ void TestCopyToDiscardIteratorZipped()
 
   thrust::host_vector<T> h_output(5);
   thrust::device_vector<T> d_output(5);
-  thrust::discard_iterator<> reference(5);
+  cuda::discard_iterator reference(5);
 
-  using Tuple1 = cuda::std::tuple<thrust::discard_iterator<>, thrust::host_vector<T>::iterator>;
-  using Tuple2 = cuda::std::tuple<thrust::discard_iterator<>, thrust::device_vector<T>::iterator>;
+  using Tuple1 = cuda::std::tuple<cuda::discard_iterator, thrust::host_vector<T>::iterator>;
+  using Tuple2 = cuda::std::tuple<cuda::discard_iterator, thrust::device_vector<T>::iterator>;
 
   using ZipIterator1 = thrust::zip_iterator<Tuple1>;
   using ZipIterator2 = thrust::zip_iterator<Tuple2>;
@@ -91,13 +90,13 @@ void TestCopyToDiscardIteratorZipped()
   ZipIterator1 h_result = thrust::copy(
     thrust::make_zip_iterator(h_input.begin(), h_input.begin()),
     thrust::make_zip_iterator(h_input.end(), h_input.end()),
-    thrust::make_zip_iterator(thrust::make_discard_iterator(), h_output.begin()));
+    thrust::make_zip_iterator(cuda::make_discard_iterator(), h_output.begin()));
 
   // copy from device_vector
   ZipIterator2 d_result = thrust::copy(
     thrust::make_zip_iterator(d_input.begin(), d_input.begin()),
     thrust::make_zip_iterator(d_input.end(), d_input.end()),
-    thrust::make_zip_iterator(thrust::make_discard_iterator(), d_output.begin()));
+    thrust::make_zip_iterator(cuda::make_discard_iterator(), d_output.begin()));
 
   ASSERT_EQUAL(h_output, h_input);
   ASSERT_EQUAL(d_output, d_input);
diff --git a/thrust/testing/copy_n.cu b/thrust/testing/copy_n.cu
index 9ee3039a706..b90e7e7d368 100644
--- a/thrust/testing/copy_n.cu
+++ b/thrust/testing/copy_n.cu
@@ -1,6 +1,5 @@
 #include <thrust/copy.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/sequence.h>
@@ -44,14 +43,12 @@ void TestCopyNToDiscardIterator()
   thrust::device_vector<T> d_input = h_input;
 
   // copy from host_vector
-  thrust::discard_iterator<> h_result =
-    thrust::copy_n(h_input.begin(), h_input.size(), thrust::make_discard_iterator());
+  cuda::discard_iterator h_result = thrust::copy_n(h_input.begin(), h_input.size(), cuda::make_discard_iterator());
 
   // copy from device_vector
-  thrust::discard_iterator<> d_result =
-    thrust::copy_n(d_input.begin(), d_input.size(), thrust::make_discard_iterator());
+  cuda::discard_iterator d_result = thrust::copy_n(d_input.begin(), d_input.size(), cuda::make_discard_iterator());
 
-  thrust::discard_iterator<> reference(5);
+  cuda::discard_iterator reference(5);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/cuda/partition.cu b/thrust/testing/cuda/partition.cu
index 3fdd51af390..9d14b7646d3 100644
--- a/thrust/testing/cuda/partition.cu
+++ b/thrust/testing/cuda/partition.cu
@@ -1,8 +1,9 @@
 #include <thrust/count.h>
 #include <thrust/execution_policy.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/partition.h>
 
+#include <cuda/iterator>
+
 #include "thrust/detail/raw_pointer_cast.h"
 #include <unittest/unittest.h>
 
@@ -634,7 +635,7 @@ void TestPartitionIfWithMagnitude(int magnitude)
     constexpr bool negate_matches = false;
     auto select_op                = mod_n<offset_t>{match_every_nth, negate_matches};
     auto partitioned_out_ends =
-      thrust::stable_partition_copy(begin, end, partitioned_out.begin(), thrust::make_discard_iterator(), select_op);
+      thrust::stable_partition_copy(begin, end, partitioned_out.begin(), cuda::make_discard_iterator(), select_op);
     const auto selected_out_end = partitioned_out_ends.first;
 
     // Ensure number of selected items are correct
@@ -658,7 +659,7 @@ void TestPartitionIfWithMagnitude(int magnitude)
     constexpr bool negate_matches = true;
     auto select_op                = mod_n<offset_t>{match_every_nth, negate_matches};
     const auto partitioned_out_ends =
-      thrust::stable_partition_copy(begin, end, thrust::make_discard_iterator(), partitioned_out.begin(), select_op);
+      thrust::stable_partition_copy(begin, end, cuda::make_discard_iterator(), partitioned_out.begin(), select_op);
     const auto rejected_out_end = partitioned_out_ends.second;
 
     // Ensure number of rejected items are correct
diff --git a/thrust/testing/cuda/set_intersection.cu b/thrust/testing/cuda/set_intersection.cu
index 11239c74eae..d8c459e3aa8 100644
--- a/thrust/testing/cuda/set_intersection.cu
+++ b/thrust/testing/cuda/set_intersection.cu
@@ -1,9 +1,10 @@
 #include <thrust/extrema.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/set_operations.h>
 #include <thrust/sort.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 #ifdef THRUST_TEST_DEVICE_SIDE
diff --git a/thrust/testing/discard_iterator.cu b/thrust/testing/discard_iterator.cu
index bf39ec1d9d3..a5aa5cb8444 100644
--- a/thrust/testing/discard_iterator.cu
+++ b/thrust/testing/discard_iterator.cu
@@ -1,3 +1,5 @@
+#define CCCL_IGNORE_DEPRECATED_API
+
 #include <thrust/iterator/discard_iterator.h>
 
 #include <cuda/std/tuple>
diff --git a/thrust/testing/fill.cu b/thrust/testing/fill.cu
index 923758eb219..f1464d7522c 100644
--- a/thrust/testing/fill.cu
+++ b/thrust/testing/fill.cu
@@ -1,8 +1,9 @@
 #include <thrust/fill.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/zip_iterator.h>
 
+#include <cuda/iterator>
+
 #include <algorithm>
 
 #include <unittest/unittest.h>
@@ -42,11 +43,9 @@ DECLARE_VECTOR_UNITTEST(TestFillSimple);
 void TestFillDiscardIterator()
 {
   // there's no result to check because fill returns void
-  thrust::fill(
-    thrust::discard_iterator<thrust::host_system_tag>(), thrust::discard_iterator<thrust::host_system_tag>(10), 13);
+  thrust::fill(thrust::host, cuda::discard_iterator(), cuda::discard_iterator(10), 13);
 
-  thrust::fill(
-    thrust::discard_iterator<thrust::device_system_tag>(), thrust::discard_iterator<thrust::device_system_tag>(10), 13);
+  thrust::fill(thrust::device, cuda::discard_iterator(), cuda::discard_iterator(10), 13);
 }
 DECLARE_UNITTEST(TestFillDiscardIterator);
 
@@ -139,13 +138,11 @@ DECLARE_VECTOR_UNITTEST(TestFillNSimple);
 
 void TestFillNDiscardIterator()
 {
-  thrust::discard_iterator<thrust::host_system_tag> h_result =
-    thrust::fill_n(thrust::discard_iterator<thrust::host_system_tag>(), 10, 13);
+  cuda::discard_iterator h_result = thrust::fill_n(thrust::host, cuda::discard_iterator(), 10, 13);
 
-  thrust::discard_iterator<thrust::device_system_tag> d_result =
-    thrust::fill_n(thrust::discard_iterator<thrust::device_system_tag>(), 10, 13);
+  cuda::discard_iterator d_result = thrust::fill_n(thrust::device, cuda::discard_iterator(), 10, 13);
 
-  thrust::discard_iterator<> reference(10);
+  cuda::discard_iterator reference(10);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/gather.cu b/thrust/testing/gather.cu
index 7ea006c441d..b95e55d8341 100644
--- a/thrust/testing/gather.cu
+++ b/thrust/testing/gather.cu
@@ -1,10 +1,11 @@
 #include <thrust/fill.h>
 #include <thrust/gather.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/sequence.h>
 
+#include <cuda/iterator>
+
 #include <algorithm>
 
 #include <unittest/unittest.h>
@@ -121,13 +122,13 @@ void TestGatherToDiscardIterator(const size_t n)
 
   thrust::device_vector<unsigned int> d_map = h_map;
 
-  thrust::discard_iterator<> h_result =
-    thrust::gather(h_map.begin(), h_map.end(), h_source.begin(), thrust::make_discard_iterator());
+  cuda::discard_iterator h_result =
+    thrust::gather(h_map.begin(), h_map.end(), h_source.begin(), cuda::make_discard_iterator());
 
-  thrust::discard_iterator<> d_result =
-    thrust::gather(d_map.begin(), d_map.end(), d_source.begin(), thrust::make_discard_iterator());
+  cuda::discard_iterator d_result =
+    thrust::gather(d_map.begin(), d_map.end(), d_source.begin(), cuda::make_discard_iterator());
 
-  thrust::discard_iterator<> reference(n);
+  cuda::discard_iterator reference(n);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
@@ -291,23 +292,23 @@ void TestGatherIfToDiscardIterator(const size_t n)
 
   thrust::device_vector<unsigned int> d_stencil = h_stencil;
 
-  thrust::discard_iterator<> h_result = thrust::gather_if(
+  cuda::discard_iterator h_result = thrust::gather_if(
     h_map.begin(),
     h_map.end(),
     h_stencil.begin(),
     h_source.begin(),
-    thrust::make_discard_iterator(),
+    cuda::make_discard_iterator(),
     is_even_gather_if<unsigned int>());
 
-  thrust::discard_iterator<> d_result = thrust::gather_if(
+  cuda::discard_iterator d_result = thrust::gather_if(
     d_map.begin(),
     d_map.end(),
     d_stencil.begin(),
     d_source.begin(),
-    thrust::make_discard_iterator(),
+    cuda::make_discard_iterator(),
     is_even_gather_if<unsigned int>());
 
-  thrust::discard_iterator<> reference(n);
+  cuda::discard_iterator reference(n);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/generate.cu b/thrust/testing/generate.cu
index e949110c14a..512af5adcae 100644
--- a/thrust/testing/generate.cu
+++ b/thrust/testing/generate.cu
@@ -1,7 +1,8 @@
 #include <thrust/generate.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 _CCCL_DIAG_PUSH
@@ -96,11 +97,11 @@ void TestGenerateToDiscardIterator(const size_t)
   T value = 13;
   return_value<T> f(value);
 
-  thrust::discard_iterator<thrust::host_system_tag> h_first;
-  thrust::generate(h_first, h_first + 10, f);
+  cuda::discard_iterator h_first;
+  thrust::generate(thrust::host, h_first, h_first + 10, f);
 
-  thrust::discard_iterator<thrust::device_system_tag> d_first;
-  thrust::generate(d_first, d_first + 10, f);
+  cuda::discard_iterator d_first;
+  thrust::generate(thrust::device, d_first, d_first + 10, f);
 
   // there's nothing to actually check except that it compiles
 }
@@ -165,13 +166,11 @@ void TestGenerateNToDiscardIterator(const size_t n)
   T value = 13;
   return_value<T> f(value);
 
-  thrust::discard_iterator<thrust::host_system_tag> h_result =
-    thrust::generate_n(thrust::discard_iterator<thrust::host_system_tag>(), n, f);
+  cuda::discard_iterator h_result = thrust::generate_n(thrust::host, cuda::discard_iterator(), n, f);
 
-  thrust::discard_iterator<thrust::device_system_tag> d_result =
-    thrust::generate_n(thrust::discard_iterator<thrust::device_system_tag>(), n, f);
+  cuda::discard_iterator d_result = thrust::generate_n(thrust::device, cuda::discard_iterator(), n, f);
 
-  thrust::discard_iterator<> reference(n);
+  cuda::discard_iterator reference(n);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/merge.cu b/thrust/testing/merge.cu
index 1720c9ec2f1..0f8cf87e113 100644
--- a/thrust/testing/merge.cu
+++ b/thrust/testing/merge.cu
@@ -1,11 +1,12 @@
 #include <thrust/extrema.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/merge.h>
 #include <thrust/sort.h>
 #include <thrust/unique.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename Vector>
@@ -110,10 +111,10 @@ void TestMergeToDiscardIterator(size_t n)
   const thrust::device_vector<T> d_a = h_a;
   const thrust::device_vector<T> d_b = h_b;
 
-  const auto h_result = thrust::merge(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), thrust::make_discard_iterator());
-  const auto d_result = thrust::merge(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), thrust::make_discard_iterator());
+  const auto h_result = thrust::merge(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), cuda::make_discard_iterator());
+  const auto d_result = thrust::merge(d_a.begin(), d_a.end(), d_b.begin(), d_b.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);
diff --git a/thrust/testing/merge_by_key.cu b/thrust/testing/merge_by_key.cu
index 28a09697d21..66ef1b9ddab 100644
--- a/thrust/testing/merge_by_key.cu
+++ b/thrust/testing/merge_by_key.cu
@@ -1,5 +1,4 @@
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/merge.h>
@@ -219,7 +218,7 @@ void TestMergeByKeyToDiscardIterator(size_t n)
   const thrust::device_vector<T> d_a_vals = h_a_vals;
   const thrust::device_vector<T> d_b_vals = h_b_vals;
 
-  using discard_pair = cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>>;
+  using discard_pair = cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator>;
 
   const discard_pair h_result = thrust::merge_by_key(
     h_a_keys.begin(),
@@ -228,8 +227,8 @@ void TestMergeByKeyToDiscardIterator(size_t n)
     h_b_keys.end(),
     h_a_vals.begin(),
     h_b_vals.begin(),
-    thrust::make_discard_iterator(),
-    thrust::make_discard_iterator());
+    cuda::make_discard_iterator(),
+    cuda::make_discard_iterator());
 
   const discard_pair d_result = thrust::merge_by_key(
     d_a_keys.begin(),
@@ -238,10 +237,10 @@ void TestMergeByKeyToDiscardIterator(size_t n)
     d_b_keys.end(),
     d_a_vals.begin(),
     d_b_vals.begin(),
-    thrust::make_discard_iterator(),
-    thrust::make_discard_iterator());
+    cuda::make_discard_iterator(),
+    cuda::make_discard_iterator());
 
-  const thrust::discard_iterator<> reference(2 * n);
+  const cuda::discard_iterator reference(2 * n);
 
   ASSERT_EQUAL_QUIET(reference, h_result.first);
   ASSERT_EQUAL_QUIET(reference, h_result.second);
@@ -307,7 +306,7 @@ void TestMergeByKeyFromCuDFDremel()
     thrust::make_counting_iterator(curr_rep_values_size),
     input_parent_zip_it,
     input_child_zip_it,
-    thrust::make_discard_iterator(),
+    cuda::make_discard_iterator(),
     output_zip_it);
 
   thrust::device_vector<std::uint8_t> reference_rep_level(max_vals_size);
diff --git a/thrust/testing/partition.cu b/thrust/testing/partition.cu
index 4d472b9ac9d..6fb02809a19 100644
--- a/thrust/testing/partition.cu
+++ b/thrust/testing/partition.cu
@@ -1,10 +1,11 @@
 #include <thrust/count.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/partition.h>
 #include <thrust/sort.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 #if _CCCL_COMPILER(GCC, >=, 11) && _CCCL_COMPILER(GCC, <, 12)
@@ -408,14 +409,14 @@ struct TestPartitionCopyToDiscardIterator
     std::ptrdiff_t n_false = n - n_true;
 
     // mask both ranges
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> h_result1 = thrust::partition_copy(
-      h_data.begin(), h_data.end(), thrust::make_discard_iterator(), thrust::make_discard_iterator(), is_even<T>());
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> h_result1 = thrust::partition_copy(
+      h_data.begin(), h_data.end(), cuda::make_discard_iterator(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> d_result1 = thrust::partition_copy(
-      d_data.begin(), d_data.end(), thrust::make_discard_iterator(), thrust::make_discard_iterator(), is_even<T>());
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> d_result1 = thrust::partition_copy(
+      d_data.begin(), d_data.end(), cuda::make_discard_iterator(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> reference1 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), thrust::make_discard_iterator(n_false));
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> reference1 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), cuda::make_discard_iterator(n_false));
 
     ASSERT_EQUAL_QUIET(reference1, h_result1);
     ASSERT_EQUAL_QUIET(reference1, d_result1);
@@ -424,19 +425,17 @@ struct TestPartitionCopyToDiscardIterator
     thrust::host_vector<T> h_trues(n_true);
     thrust::device_vector<T> d_trues(n_true);
 
-    cuda::std::pair<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>> h_result2 =
-      thrust::partition_copy(
-        h_data.begin(), h_data.end(), h_trues.begin(), thrust::make_discard_iterator(), is_even<T>());
+    cuda::std::pair<typename thrust::host_vector<T>::iterator, cuda::discard_iterator> h_result2 =
+      thrust::partition_copy(h_data.begin(), h_data.end(), h_trues.begin(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>> d_result2 =
-      thrust::partition_copy(
-        d_data.begin(), d_data.end(), d_trues.begin(), thrust::make_discard_iterator(), is_even<T>());
+    cuda::std::pair<typename thrust::device_vector<T>::iterator, cuda::discard_iterator> d_result2 =
+      thrust::partition_copy(d_data.begin(), d_data.end(), d_trues.begin(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>> h_reference2 =
-      cuda::std::make_pair(h_trues.begin() + n_true, thrust::make_discard_iterator(n_false));
+    cuda::std::pair<typename thrust::host_vector<T>::iterator, cuda::discard_iterator> h_reference2 =
+      cuda::std::make_pair(h_trues.begin() + n_true, cuda::make_discard_iterator(n_false));
 
-    cuda::std::pair<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>> d_reference2 =
-      cuda::std::make_pair(d_trues.begin() + n_true, thrust::make_discard_iterator(n_false));
+    cuda::std::pair<typename thrust::device_vector<T>::iterator, cuda::discard_iterator> d_reference2 =
+      cuda::std::make_pair(d_trues.begin() + n_true, cuda::make_discard_iterator(n_false));
 
     ASSERT_EQUAL(h_trues, d_trues);
     ASSERT_EQUAL_QUIET(h_reference2, h_result2);
@@ -446,19 +445,19 @@ struct TestPartitionCopyToDiscardIterator
     thrust::host_vector<T> h_falses(n_false);
     thrust::device_vector<T> d_falses(n_false);
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<T>::iterator> h_result3 =
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<T>::iterator> h_result3 =
       thrust::partition_copy(
-        h_data.begin(), h_data.end(), thrust::make_discard_iterator(), h_falses.begin(), is_even<T>());
+        h_data.begin(), h_data.end(), cuda::make_discard_iterator(), h_falses.begin(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<T>::iterator> d_result3 =
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<T>::iterator> d_result3 =
       thrust::partition_copy(
-        d_data.begin(), d_data.end(), thrust::make_discard_iterator(), d_falses.begin(), is_even<T>());
+        d_data.begin(), d_data.end(), cuda::make_discard_iterator(), d_falses.begin(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<T>::iterator> h_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), h_falses.begin() + n_false);
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<T>::iterator> h_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), h_falses.begin() + n_false);
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<T>::iterator> d_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), d_falses.begin() + n_false);
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<T>::iterator> d_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), d_falses.begin() + n_false);
 
     ASSERT_EQUAL(h_falses, d_falses);
     ASSERT_EQUAL_QUIET(h_reference3, h_result3);
@@ -482,24 +481,24 @@ struct TestPartitionCopyStencilToDiscardIterator
     std::ptrdiff_t n_false = n - n_true;
 
     // mask both ranges
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> h_result1 = thrust::partition_copy(
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> h_result1 = thrust::partition_copy(
       h_data.begin(),
       h_data.end(),
       h_stencil.begin(),
-      thrust::make_discard_iterator(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> d_result1 = thrust::partition_copy(
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> d_result1 = thrust::partition_copy(
       d_data.begin(),
       d_data.end(),
       d_stencil.begin(),
-      thrust::make_discard_iterator(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> reference1 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), thrust::make_discard_iterator(n_false));
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> reference1 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), cuda::make_discard_iterator(n_false));
 
     ASSERT_EQUAL_QUIET(reference1, h_result1);
     ASSERT_EQUAL_QUIET(reference1, d_result1);
@@ -508,19 +507,19 @@ struct TestPartitionCopyStencilToDiscardIterator
     thrust::host_vector<T> h_trues(n_true);
     thrust::device_vector<T> d_trues(n_true);
 
-    cuda::std::pair<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>> h_result2 =
+    cuda::std::pair<typename thrust::host_vector<T>::iterator, cuda::discard_iterator> h_result2 =
       thrust::partition_copy(
-        h_data.begin(), h_data.end(), h_stencil.begin(), h_trues.begin(), thrust::make_discard_iterator(), is_even<T>());
+        h_data.begin(), h_data.end(), h_stencil.begin(), h_trues.begin(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>> d_result2 =
+    cuda::std::pair<typename thrust::device_vector<T>::iterator, cuda::discard_iterator> d_result2 =
       thrust::partition_copy(
-        d_data.begin(), d_data.end(), d_stencil.begin(), d_trues.begin(), thrust::make_discard_iterator(), is_even<T>());
+        d_data.begin(), d_data.end(), d_stencil.begin(), d_trues.begin(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>> h_reference2 =
-      cuda::std::make_pair(h_trues.begin() + n_true, thrust::make_discard_iterator(n_false));
+    cuda::std::pair<typename thrust::host_vector<T>::iterator, cuda::discard_iterator> h_reference2 =
+      cuda::std::make_pair(h_trues.begin() + n_true, cuda::make_discard_iterator(n_false));
 
-    cuda::std::pair<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>> d_reference2 =
-      cuda::std::make_pair(d_trues.begin() + n_true, thrust::make_discard_iterator(n_false));
+    cuda::std::pair<typename thrust::device_vector<T>::iterator, cuda::discard_iterator> d_reference2 =
+      cuda::std::make_pair(d_trues.begin() + n_true, cuda::make_discard_iterator(n_false));
 
     ASSERT_EQUAL(h_trues, d_trues);
     ASSERT_EQUAL_QUIET(h_reference2, h_result2);
@@ -530,29 +529,19 @@ struct TestPartitionCopyStencilToDiscardIterator
     thrust::host_vector<T> h_falses(n_false);
     thrust::device_vector<T> d_falses(n_false);
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<T>::iterator> h_result3 =
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<T>::iterator> h_result3 =
       thrust::partition_copy(
-        h_data.begin(),
-        h_data.end(),
-        h_stencil.begin(),
-        thrust::make_discard_iterator(),
-        h_falses.begin(),
-        is_even<T>());
-
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<T>::iterator> d_result3 =
+        h_data.begin(), h_data.end(), h_stencil.begin(), cuda::make_discard_iterator(), h_falses.begin(), is_even<T>());
+
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<T>::iterator> d_result3 =
       thrust::partition_copy(
-        d_data.begin(),
-        d_data.end(),
-        d_stencil.begin(),
-        thrust::make_discard_iterator(),
-        d_falses.begin(),
-        is_even<T>());
+        d_data.begin(), d_data.end(), d_stencil.begin(), cuda::make_discard_iterator(), d_falses.begin(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<T>::iterator> h_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), h_falses.begin() + n_false);
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<T>::iterator> h_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), h_falses.begin() + n_false);
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<T>::iterator> d_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), d_falses.begin() + n_false);
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<T>::iterator> d_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), d_falses.begin() + n_false);
 
     ASSERT_EQUAL(h_falses, d_falses);
     ASSERT_EQUAL_QUIET(h_reference3, h_result3);
@@ -666,14 +655,14 @@ struct TestStablePartitionCopyToDiscardIterator
     std::ptrdiff_t n_false = n - n_true;
 
     // mask both ranges
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> h_result1 = thrust::stable_partition_copy(
-      h_data.begin(), h_data.end(), thrust::make_discard_iterator(), thrust::make_discard_iterator(), is_even<T>());
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> h_result1 = thrust::stable_partition_copy(
+      h_data.begin(), h_data.end(), cuda::make_discard_iterator(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> d_result1 = thrust::stable_partition_copy(
-      d_data.begin(), d_data.end(), thrust::make_discard_iterator(), thrust::make_discard_iterator(), is_even<T>());
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> d_result1 = thrust::stable_partition_copy(
+      d_data.begin(), d_data.end(), cuda::make_discard_iterator(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> reference1 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), thrust::make_discard_iterator(n_false));
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> reference1 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), cuda::make_discard_iterator(n_false));
 
     ASSERT_EQUAL_QUIET(reference1, h_result1);
     ASSERT_EQUAL_QUIET(reference1, d_result1);
@@ -682,19 +671,19 @@ struct TestStablePartitionCopyToDiscardIterator
     thrust::host_vector<T> h_trues(n_true);
     thrust::device_vector<T> d_trues(n_true);
 
-    cuda::std::pair<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>> h_result2 =
+    cuda::std::pair<typename thrust::host_vector<T>::iterator, cuda::discard_iterator> h_result2 =
       thrust::stable_partition_copy(
-        h_data.begin(), h_data.end(), h_trues.begin(), thrust::make_discard_iterator(), is_even<T>());
+        h_data.begin(), h_data.end(), h_trues.begin(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>> d_result2 =
+    cuda::std::pair<typename thrust::device_vector<T>::iterator, cuda::discard_iterator> d_result2 =
       thrust::stable_partition_copy(
-        d_data.begin(), d_data.end(), d_trues.begin(), thrust::make_discard_iterator(), is_even<T>());
+        d_data.begin(), d_data.end(), d_trues.begin(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>> h_reference2 =
-      cuda::std::make_pair(h_trues.begin() + n_true, thrust::make_discard_iterator(n_false));
+    cuda::std::pair<typename thrust::host_vector<T>::iterator, cuda::discard_iterator> h_reference2 =
+      cuda::std::make_pair(h_trues.begin() + n_true, cuda::make_discard_iterator(n_false));
 
-    cuda::std::pair<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>> d_reference2 =
-      cuda::std::make_pair(d_trues.begin() + n_true, thrust::make_discard_iterator(n_false));
+    cuda::std::pair<typename thrust::device_vector<T>::iterator, cuda::discard_iterator> d_reference2 =
+      cuda::std::make_pair(d_trues.begin() + n_true, cuda::make_discard_iterator(n_false));
 
     ASSERT_EQUAL(h_trues, d_trues);
     ASSERT_EQUAL_QUIET(h_reference2, h_result2);
@@ -704,19 +693,19 @@ struct TestStablePartitionCopyToDiscardIterator
     thrust::host_vector<T> h_falses(n_false);
     thrust::device_vector<T> d_falses(n_false);
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<T>::iterator> h_result3 =
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<T>::iterator> h_result3 =
       thrust::stable_partition_copy(
-        h_data.begin(), h_data.end(), thrust::make_discard_iterator(), h_falses.begin(), is_even<T>());
+        h_data.begin(), h_data.end(), cuda::make_discard_iterator(), h_falses.begin(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<T>::iterator> d_result3 =
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<T>::iterator> d_result3 =
       thrust::stable_partition_copy(
-        d_data.begin(), d_data.end(), thrust::make_discard_iterator(), d_falses.begin(), is_even<T>());
+        d_data.begin(), d_data.end(), cuda::make_discard_iterator(), d_falses.begin(), is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<T>::iterator> h_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), h_falses.begin() + n_false);
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<T>::iterator> h_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), h_falses.begin() + n_false);
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<T>::iterator> d_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), d_falses.begin() + n_false);
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<T>::iterator> d_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), d_falses.begin() + n_false);
 
     ASSERT_EQUAL(h_falses, d_falses);
     ASSERT_EQUAL_QUIET(h_reference3, h_result3);
@@ -741,24 +730,24 @@ struct TestStablePartitionCopyStencilToDiscardIterator
     std::ptrdiff_t n_false = n - n_true;
 
     // mask both ranges
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> h_result1 = thrust::stable_partition_copy(
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> h_result1 = thrust::stable_partition_copy(
       h_data.begin(),
       h_data.end(),
       h_stencil.begin(),
-      thrust::make_discard_iterator(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> d_result1 = thrust::stable_partition_copy(
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> d_result1 = thrust::stable_partition_copy(
       d_data.begin(),
       d_data.end(),
       d_stencil.begin(),
-      thrust::make_discard_iterator(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       is_even<T>());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> reference1 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), thrust::make_discard_iterator(n_false));
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> reference1 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), cuda::make_discard_iterator(n_false));
 
     ASSERT_EQUAL_QUIET(reference1, h_result1);
     ASSERT_EQUAL_QUIET(reference1, d_result1);
@@ -767,19 +756,19 @@ struct TestStablePartitionCopyStencilToDiscardIterator
     thrust::host_vector<T> h_trues(n_true);
     thrust::device_vector<T> d_trues(n_true);
 
-    cuda::std::pair<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>> h_result2 =
+    cuda::std::pair<typename thrust::host_vector<T>::iterator, cuda::discard_iterator> h_result2 =
       thrust::stable_partition_copy(
-        h_data.begin(), h_data.end(), h_stencil.begin(), h_trues.begin(), thrust::make_discard_iterator(), is_even<T>());
+        h_data.begin(), h_data.end(), h_stencil.begin(), h_trues.begin(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>> d_result2 =
+    cuda::std::pair<typename thrust::device_vector<T>::iterator, cuda::discard_iterator> d_result2 =
       thrust::stable_partition_copy(
-        d_data.begin(), d_data.end(), d_stencil.begin(), d_trues.begin(), thrust::make_discard_iterator(), is_even<T>());
+        d_data.begin(), d_data.end(), d_stencil.begin(), d_trues.begin(), cuda::make_discard_iterator(), is_even<T>());
 
-    cuda::std::pair<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>> h_reference2 =
-      cuda::std::make_pair(h_trues.begin() + n_true, thrust::make_discard_iterator(n_false));
+    cuda::std::pair<typename thrust::host_vector<T>::iterator, cuda::discard_iterator> h_reference2 =
+      cuda::std::make_pair(h_trues.begin() + n_true, cuda::make_discard_iterator(n_false));
 
-    cuda::std::pair<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>> d_reference2 =
-      cuda::std::make_pair(d_trues.begin() + n_true, thrust::make_discard_iterator(n_false));
+    cuda::std::pair<typename thrust::device_vector<T>::iterator, cuda::discard_iterator> d_reference2 =
+      cuda::std::make_pair(d_trues.begin() + n_true, cuda::make_discard_iterator(n_false));
 
     ASSERT_EQUAL(h_trues, d_trues);
     ASSERT_EQUAL_QUIET(h_reference2, h_result2);
@@ -789,29 +778,19 @@ struct TestStablePartitionCopyStencilToDiscardIterator
     thrust::host_vector<T> h_falses(n_false);
     thrust::device_vector<T> d_falses(n_false);
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<T>::iterator> h_result3 =
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<T>::iterator> h_result3 =
       thrust::stable_partition_copy(
-        h_data.begin(),
-        h_data.end(),
-        h_stencil.begin(),
-        thrust::make_discard_iterator(),
-        h_falses.begin(),
-        is_even<T>());
-
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<T>::iterator> d_result3 =
+        h_data.begin(), h_data.end(), h_stencil.begin(), cuda::make_discard_iterator(), h_falses.begin(), is_even<T>());
+
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<T>::iterator> d_result3 =
       thrust::stable_partition_copy(
-        d_data.begin(),
-        d_data.end(),
-        d_stencil.begin(),
-        thrust::make_discard_iterator(),
-        d_falses.begin(),
-        is_even<T>());
-
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<T>::iterator> h_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), h_falses.begin() + n_false);
-
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<T>::iterator> d_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(n_true), d_falses.begin() + n_false);
+        d_data.begin(), d_data.end(), d_stencil.begin(), cuda::make_discard_iterator(), d_falses.begin(), is_even<T>());
+
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<T>::iterator> h_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), h_falses.begin() + n_false);
+
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<T>::iterator> d_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(n_true), d_falses.begin() + n_false);
 
     ASSERT_EQUAL(h_falses, d_falses);
     ASSERT_EQUAL_QUIET(h_reference3, h_result3);
diff --git a/thrust/testing/raw_reference_cast.cu b/thrust/testing/raw_reference_cast.cu
index 947f2847c26..0471b838a91 100644
--- a/thrust/testing/raw_reference_cast.cu
+++ b/thrust/testing/raw_reference_cast.cu
@@ -1,8 +1,9 @@
 #include <thrust/detail/raw_reference_cast.h>
 #include <thrust/device_vector.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 
+#include <cuda/iterator>
+
 #include <vector>
 
 #include <unittest/unittest.h>
diff --git a/thrust/testing/reduce_by_key.cu b/thrust/testing/reduce_by_key.cu
index 2c0133b7900..282e2493a58 100644
--- a/thrust/testing/reduce_by_key.cu
+++ b/thrust/testing/reduce_by_key.cu
@@ -1,8 +1,9 @@
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/reduce.h>
 #include <thrust/unique.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename T>
@@ -174,13 +175,13 @@ struct TestReduceByKeyToDiscardIterator
     // discard key output
     size_t h_size =
       thrust::reduce_by_key(
-        h_keys.begin(), h_keys.end(), h_vals.begin(), thrust::make_discard_iterator(), h_vals_output.begin())
+        h_keys.begin(), h_keys.end(), h_vals.begin(), cuda::make_discard_iterator(), h_vals_output.begin())
         .second
       - h_vals_output.begin();
 
     size_t d_size =
       thrust::reduce_by_key(
-        d_keys.begin(), d_keys.end(), d_vals.begin(), thrust::make_discard_iterator(), d_vals_output.begin())
+        d_keys.begin(), d_keys.end(), d_vals.begin(), cuda::make_discard_iterator(), d_vals_output.begin())
         .second
       - d_vals_output.begin();
 
diff --git a/thrust/testing/remove.cu b/thrust/testing/remove.cu
index b2c61e19ee6..4bfe59ddb86 100644
--- a/thrust/testing/remove.cu
+++ b/thrust/testing/remove.cu
@@ -1,10 +1,11 @@
 #include <thrust/count.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/remove.h>
 
+#include <cuda/iterator>
+
 #include <stdexcept>
 
 #include <unittest/unittest.h>
@@ -440,13 +441,13 @@ void TestRemoveCopyToDiscardIterator(const size_t n)
   size_t num_zeros    = thrust::count(h_data.begin(), h_data.end(), T(0));
   size_t num_nonzeros = h_data.size() - num_zeros;
 
-  thrust::discard_iterator<> h_result =
-    thrust::remove_copy(h_data.begin(), h_data.end(), thrust::make_discard_iterator(), T(0));
+  cuda::discard_iterator h_result =
+    thrust::remove_copy(h_data.begin(), h_data.end(), cuda::make_discard_iterator(), T(0));
 
-  thrust::discard_iterator<> d_result =
-    thrust::remove_copy(d_data.begin(), d_data.end(), thrust::make_discard_iterator(), T(0));
+  cuda::discard_iterator d_result =
+    thrust::remove_copy(d_data.begin(), d_data.end(), cuda::make_discard_iterator(), T(0));
 
-  thrust::discard_iterator<> reference(num_nonzeros);
+  cuda::discard_iterator reference(num_nonzeros);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
@@ -465,8 +466,8 @@ void TestRemoveCopyToDiscardIteratorZipped(const size_t n)
   size_t num_zeros    = thrust::count(h_data.begin(), h_data.end(), T(0));
   size_t num_nonzeros = h_data.size() - num_zeros;
 
-  using Tuple1 = cuda::std::tuple<typename thrust::host_vector<T>::iterator, thrust::discard_iterator<>>;
-  using Tuple2 = cuda::std::tuple<typename thrust::device_vector<T>::iterator, thrust::discard_iterator<>>;
+  using Tuple1 = cuda::std::tuple<typename thrust::host_vector<T>::iterator, cuda::discard_iterator>;
+  using Tuple2 = cuda::std::tuple<typename thrust::device_vector<T>::iterator, cuda::discard_iterator>;
 
   using ZipIterator1 = thrust::zip_iterator<Tuple1>;
   using ZipIterator2 = thrust::zip_iterator<Tuple2>;
@@ -474,16 +475,16 @@ void TestRemoveCopyToDiscardIteratorZipped(const size_t n)
   ZipIterator1 h_result = thrust::remove_copy(
     thrust::make_zip_iterator(h_data.begin(), h_data.begin()),
     thrust::make_zip_iterator(h_data.end(), h_data.end()),
-    thrust::make_zip_iterator(h_output.begin(), thrust::make_discard_iterator()),
+    thrust::make_zip_iterator(h_output.begin(), cuda::make_discard_iterator()),
     cuda::std::tuple(T(0), T(0)));
 
   ZipIterator2 d_result = thrust::remove_copy(
     thrust::make_zip_iterator(d_data.begin(), d_data.begin()),
     thrust::make_zip_iterator(d_data.end(), d_data.end()),
-    thrust::make_zip_iterator(d_output.begin(), thrust::make_discard_iterator()),
+    thrust::make_zip_iterator(d_output.begin(), cuda::make_discard_iterator()),
     cuda::std::tuple(T(0), T(0)));
 
-  thrust::discard_iterator<> reference(num_nonzeros);
+  cuda::discard_iterator reference(num_nonzeros);
 
   ASSERT_EQUAL(h_output, d_output);
   ASSERT_EQUAL_QUIET(reference, cuda::std::get<1>(h_result.get_iterator_tuple()));
@@ -522,13 +523,13 @@ void TestRemoveCopyIfToDiscardIterator(const size_t n)
 
   size_t num_false = thrust::count_if(h_data.begin(), h_data.end(), ::cuda::std::not_fn(is_true<T>()));
 
-  thrust::discard_iterator<> h_result =
-    thrust::remove_copy_if(h_data.begin(), h_data.end(), thrust::make_discard_iterator(), is_true<T>());
+  cuda::discard_iterator h_result =
+    thrust::remove_copy_if(h_data.begin(), h_data.end(), cuda::make_discard_iterator(), is_true<T>());
 
-  thrust::discard_iterator<> d_result =
-    thrust::remove_copy_if(d_data.begin(), d_data.end(), thrust::make_discard_iterator(), is_true<T>());
+  cuda::discard_iterator d_result =
+    thrust::remove_copy_if(d_data.begin(), d_data.end(), cuda::make_discard_iterator(), is_true<T>());
 
-  thrust::discard_iterator<> reference(num_false);
+  cuda::discard_iterator reference(num_false);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
@@ -574,13 +575,13 @@ void TestRemoveCopyIfStencilToDiscardIterator(const size_t n)
 
   size_t num_false = thrust::count_if(h_stencil.begin(), h_stencil.end(), ::cuda::std::not_fn(is_true<T>()));
 
-  thrust::discard_iterator<> h_result = thrust::remove_copy_if(
-    h_data.begin(), h_data.end(), h_stencil.begin(), thrust::make_discard_iterator(), is_true<T>());
+  cuda::discard_iterator h_result = thrust::remove_copy_if(
+    h_data.begin(), h_data.end(), h_stencil.begin(), cuda::make_discard_iterator(), is_true<T>());
 
-  thrust::discard_iterator<> d_result = thrust::remove_copy_if(
-    d_data.begin(), d_data.end(), d_stencil.begin(), thrust::make_discard_iterator(), is_true<T>());
+  cuda::discard_iterator d_result = thrust::remove_copy_if(
+    d_data.begin(), d_data.end(), d_stencil.begin(), cuda::make_discard_iterator(), is_true<T>());
 
-  thrust::discard_iterator<> reference(num_false);
+  cuda::discard_iterator reference(num_false);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/replace.cu b/thrust/testing/replace.cu
index 34713fd2aa2..dccbb75838d 100644
--- a/thrust/testing/replace.cu
+++ b/thrust/testing/replace.cu
@@ -1,7 +1,8 @@
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/replace.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 // There is a unfortunate miscompilation of the gcc-11 vectorizer leading to OOB writes
@@ -166,13 +167,13 @@ void TestReplaceCopyToDiscardIterator(const size_t n)
   T old_value = 0;
   T new_value = 1;
 
-  thrust::discard_iterator<> h_result =
-    thrust::replace_copy(h_data.begin(), h_data.end(), thrust::make_discard_iterator(), old_value, new_value);
+  cuda::discard_iterator h_result =
+    thrust::replace_copy(h_data.begin(), h_data.end(), cuda::make_discard_iterator(), old_value, new_value);
 
-  thrust::discard_iterator<> d_result =
-    thrust::replace_copy(d_data.begin(), d_data.end(), thrust::make_discard_iterator(), old_value, new_value);
+  cuda::discard_iterator d_result =
+    thrust::replace_copy(d_data.begin(), d_data.end(), cuda::make_discard_iterator(), old_value, new_value);
 
-  thrust::discard_iterator<> reference(n);
+  cuda::discard_iterator reference(n);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
@@ -452,13 +453,13 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestReplaceCopyIfToDiscardIterator(con
   thrust::host_vector<T> h_data   = unittest::random_samples<T>(n);
   thrust::device_vector<T> d_data = h_data;
 
-  thrust::discard_iterator<> h_result =
-    thrust::replace_copy_if(h_data.begin(), h_data.end(), thrust::make_discard_iterator(), less_than_five<T>(), T{0});
+  cuda::discard_iterator h_result =
+    thrust::replace_copy_if(h_data.begin(), h_data.end(), cuda::make_discard_iterator(), less_than_five<T>(), T{0});
 
-  thrust::discard_iterator<> d_result =
-    thrust::replace_copy_if(d_data.begin(), d_data.end(), thrust::make_discard_iterator(), less_than_five<T>(), T{0});
+  cuda::discard_iterator d_result =
+    thrust::replace_copy_if(d_data.begin(), d_data.end(), cuda::make_discard_iterator(), less_than_five<T>(), T{0});
 
-  thrust::discard_iterator<> reference(n);
+  cuda::discard_iterator reference(n);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
@@ -494,13 +495,13 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestReplaceCopyIfStencilToDiscardItera
   thrust::host_vector<T> h_stencil   = unittest::random_samples<T>(n);
   thrust::device_vector<T> d_stencil = h_stencil;
 
-  thrust::discard_iterator<> h_result = thrust::replace_copy_if(
-    h_data.begin(), h_data.end(), h_stencil.begin(), thrust::make_discard_iterator(), less_than_five<T>(), T{0});
+  cuda::discard_iterator h_result = thrust::replace_copy_if(
+    h_data.begin(), h_data.end(), h_stencil.begin(), cuda::make_discard_iterator(), less_than_five<T>(), T{0});
 
-  thrust::discard_iterator<> d_result = thrust::replace_copy_if(
-    d_data.begin(), d_data.end(), d_stencil.begin(), thrust::make_discard_iterator(), less_than_five<T>(), T{0});
+  cuda::discard_iterator d_result = thrust::replace_copy_if(
+    d_data.begin(), d_data.end(), d_stencil.begin(), cuda::make_discard_iterator(), less_than_five<T>(), T{0});
 
-  thrust::discard_iterator<> reference(n);
+  cuda::discard_iterator reference(n);
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/reverse.cu b/thrust/testing/reverse.cu
index 4cbcfbc0efd..873028fc265 100644
--- a/thrust/testing/reverse.cu
+++ b/thrust/testing/reverse.cu
@@ -1,7 +1,8 @@
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/reverse.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 using ReverseTypes = unittest::type_list<unittest::int8_t, unittest::int16_t, unittest::int32_t>;
@@ -156,13 +157,11 @@ struct TestReverseCopyToDiscardIterator
     thrust::host_vector<T> h_data   = unittest::random_integers<T>(n);
     thrust::device_vector<T> d_data = h_data;
 
-    thrust::discard_iterator<> h_result =
-      thrust::reverse_copy(h_data.begin(), h_data.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator h_result = thrust::reverse_copy(h_data.begin(), h_data.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> d_result =
-      thrust::reverse_copy(d_data.begin(), d_data.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator d_result = thrust::reverse_copy(d_data.begin(), d_data.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/scan.cu b/thrust/testing/scan.cu
index c43c88d9bd5..279d6e372a3 100644
--- a/thrust/testing/scan.cu
+++ b/thrust/testing/scan.cu
@@ -3,7 +3,6 @@
 #include <thrust/device_free.h>
 #include <thrust/device_malloc.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/scan.h>
 #include <thrust/tabulate.h>
@@ -303,22 +302,22 @@ struct TestScanWithOperatorToDiscardIterator
     thrust::host_vector<T> h_input   = unittest::random_integers<T>(n);
     thrust::device_vector<T> d_input = h_input;
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
-    thrust::discard_iterator<> h_result =
-      thrust::inclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), cuda::maximum<T>{});
+    cuda::discard_iterator h_result =
+      thrust::inclusive_scan(h_input.begin(), h_input.end(), cuda::make_discard_iterator(), cuda::maximum<T>{});
 
-    thrust::discard_iterator<> d_result =
-      thrust::inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), cuda::maximum<T>{});
+    cuda::discard_iterator d_result =
+      thrust::inclusive_scan(d_input.begin(), d_input.end(), cuda::make_discard_iterator(), cuda::maximum<T>{});
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
 
-    h_result = thrust::exclusive_scan(
-      h_input.begin(), h_input.end(), thrust::make_discard_iterator(), T(13), cuda::maximum<T>{});
+    h_result =
+      thrust::exclusive_scan(h_input.begin(), h_input.end(), cuda::make_discard_iterator(), T(13), cuda::maximum<T>{});
 
-    d_result = thrust::exclusive_scan(
-      d_input.begin(), d_input.end(), thrust::make_discard_iterator(), T(13), cuda::maximum<T>{});
+    d_result =
+      thrust::exclusive_scan(d_input.begin(), d_input.end(), cuda::make_discard_iterator(), T(13), cuda::maximum<T>{});
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
@@ -375,20 +374,20 @@ struct TestScanToDiscardIterator
     thrust::host_vector<T> h_input   = unittest::random_integers<T>(n);
     thrust::device_vector<T> d_input = h_input;
 
-    thrust::discard_iterator<> h_result =
-      thrust::inclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator h_result =
+      thrust::inclusive_scan(h_input.begin(), h_input.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> d_result =
-      thrust::inclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator d_result =
+      thrust::inclusive_scan(d_input.begin(), d_input.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
 
-    h_result = thrust::exclusive_scan(h_input.begin(), h_input.end(), thrust::make_discard_iterator(), (T) 11);
+    h_result = thrust::exclusive_scan(h_input.begin(), h_input.end(), cuda::make_discard_iterator(), (T) 11);
 
-    d_result = thrust::exclusive_scan(d_input.begin(), d_input.end(), thrust::make_discard_iterator(), (T) 11);
+    d_result = thrust::exclusive_scan(d_input.begin(), d_input.end(), cuda::make_discard_iterator(), (T) 11);
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/scan_by_key.exclusive.cu b/thrust/testing/scan_by_key.exclusive.cu
index 7ef8408fd0c..36320d0d34b 100644
--- a/thrust/testing/scan_by_key.exclusive.cu
+++ b/thrust/testing/scan_by_key.exclusive.cu
@@ -1,10 +1,11 @@
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/random.h>
 #include <thrust/scan.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename Vector>
@@ -272,7 +273,7 @@ void TestScanByKeyDiscardOutput(std::size_t n)
   }
   thrust::device_vector<T> d_vals = h_vals;
 
-  auto out = thrust::make_discard_iterator();
+  auto out = cuda::make_discard_iterator();
 
   // These are no-ops, but they should compile.
   thrust::exclusive_scan_by_key(d_keys.cbegin(), d_keys.cend(), d_vals.cbegin(), out);
diff --git a/thrust/testing/scan_by_key.inclusive.cu b/thrust/testing/scan_by_key.inclusive.cu
index f7837229b50..4d9a2ccd88a 100644
--- a/thrust/testing/scan_by_key.inclusive.cu
+++ b/thrust/testing/scan_by_key.inclusive.cu
@@ -1,10 +1,11 @@
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/random.h>
 #include <thrust/scan.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename Vector>
@@ -263,7 +264,7 @@ void TestScanByKeyDiscardOutput(std::size_t n)
   }
   thrust::device_vector<T> d_vals = h_vals;
 
-  auto out = thrust::make_discard_iterator();
+  auto out = cuda::make_discard_iterator();
 
   // These are no-ops, but they should compile.
   thrust::inclusive_scan_by_key(d_keys.cbegin(), d_keys.cend(), d_vals.cbegin(), out);
diff --git a/thrust/testing/scatter.cu b/thrust/testing/scatter.cu
index 5110a3d8a21..f1c87f343c5 100644
--- a/thrust/testing/scatter.cu
+++ b/thrust/testing/scatter.cu
@@ -1,10 +1,11 @@
 #include <thrust/fill.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/scatter.h>
 #include <thrust/sequence.h>
 
+#include <cuda/iterator>
+
 #include <algorithm>
 
 #include <unittest/unittest.h>
@@ -103,8 +104,8 @@ void TestScatterToDiscardIterator(const size_t n)
 
   thrust::device_vector<unsigned int> d_map = h_map;
 
-  thrust::scatter(h_input.begin(), h_input.end(), h_map.begin(), thrust::make_discard_iterator());
-  thrust::scatter(d_input.begin(), d_input.end(), d_map.begin(), thrust::make_discard_iterator());
+  thrust::scatter(h_input.begin(), h_input.end(), h_map.begin(), cuda::make_discard_iterator());
+  thrust::scatter(d_input.begin(), d_input.end(), d_map.begin(), cuda::make_discard_iterator());
 
   // there's nothing to check -- just make sure it compiles
 }
@@ -224,14 +225,14 @@ void TestScatterIfToDiscardIterator(const size_t n)
     h_input.end(),
     h_map.begin(),
     h_map.begin(),
-    thrust::make_discard_iterator(),
+    cuda::make_discard_iterator(),
     is_even_scatter_if<unsigned int>());
   thrust::scatter_if(
     d_input.begin(),
     d_input.end(),
     d_map.begin(),
     d_map.begin(),
-    thrust::make_discard_iterator(),
+    cuda::make_discard_iterator(),
     is_even_scatter_if<unsigned int>());
 }
 DECLARE_VARIABLE_UNITTEST(TestScatterIfToDiscardIterator);
diff --git a/thrust/testing/sequence.cu b/thrust/testing/sequence.cu
index 19789c422f8..3323f7e9b1b 100644
--- a/thrust/testing/sequence.cu
+++ b/thrust/testing/sequence.cu
@@ -1,7 +1,8 @@
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/sequence.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename ForwardIterator>
@@ -89,15 +90,9 @@ void TestSequence(size_t n)
 DECLARE_VARIABLE_UNITTEST(TestSequence);
 
 template <typename T>
-void TestSequenceToDiscardIterator(size_t n)
+void TestSequenceToDiscardIterator(size_t)
 {
-  thrust::host_vector<T> h_data(n);
-  thrust::device_vector<T> d_data(n);
-
-  thrust::sequence(thrust::discard_iterator<thrust::device_system_tag>(),
-                   thrust::discard_iterator<thrust::device_system_tag>(13),
-                   T(10),
-                   T(2));
+  thrust::sequence(thrust::device, cuda::discard_iterator(), cuda::discard_iterator(13), T(10), T(2));
 
   // nothing to check -- just make sure it compiles
 }
diff --git a/thrust/testing/set_intersection.cu b/thrust/testing/set_intersection.cu
index 7ea458b57ef..80b33a20e09 100644
--- a/thrust/testing/set_intersection.cu
+++ b/thrust/testing/set_intersection.cu
@@ -1,10 +1,11 @@
 #include <thrust/extrema.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/set_operations.h>
 #include <thrust/sort.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename InputIterator1, typename InputIterator2, typename OutputIterator>
@@ -118,19 +119,19 @@ void TestSetIntersectionToDiscardIterator(const size_t n)
   thrust::device_vector<T> d_a = h_a;
   thrust::device_vector<T> d_b = h_b;
 
-  thrust::discard_iterator<> h_result;
-  thrust::discard_iterator<> d_result;
+  cuda::discard_iterator h_result;
+  cuda::discard_iterator d_result;
 
   thrust::host_vector<T> h_reference(n);
   typename thrust::host_vector<T>::iterator h_end =
     thrust::set_intersection(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), h_reference.begin());
   h_reference.erase(h_end, h_reference.end());
 
-  h_result = thrust::set_intersection(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), thrust::make_discard_iterator());
+  h_result = thrust::set_intersection(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), cuda::make_discard_iterator());
 
-  d_result = thrust::set_intersection(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), thrust::make_discard_iterator());
+  d_result = thrust::set_intersection(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), cuda::make_discard_iterator());
 
-  thrust::discard_iterator<> reference(h_reference.size());
+  cuda::discard_iterator reference(h_reference.size());
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/set_union.cu b/thrust/testing/set_union.cu
index 63dc5bb7577..b3a93569b61 100644
--- a/thrust/testing/set_union.cu
+++ b/thrust/testing/set_union.cu
@@ -1,10 +1,11 @@
 #include <thrust/extrema.h>
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/set_operations.h>
 #include <thrust/sort.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename InputIterator1, typename InputIterator2, typename OutputIterator>
@@ -134,19 +135,19 @@ void TestSetUnionToDiscardIterator(const size_t n)
   thrust::device_vector<T> d_a = h_a;
   thrust::device_vector<T> d_b = h_b;
 
-  thrust::discard_iterator<> h_result;
-  thrust::discard_iterator<> d_result;
+  cuda::discard_iterator h_result;
+  cuda::discard_iterator d_result;
 
   thrust::host_vector<T> h_reference(2 * n);
   typename thrust::host_vector<T>::iterator h_end =
     thrust::set_union(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), h_reference.begin());
   h_reference.erase(h_end, h_reference.end());
 
-  h_result = thrust::set_union(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), thrust::make_discard_iterator());
+  h_result = thrust::set_union(h_a.begin(), h_a.end(), h_b.begin(), h_b.end(), cuda::make_discard_iterator());
 
-  d_result = thrust::set_union(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), thrust::make_discard_iterator());
+  d_result = thrust::set_union(d_a.begin(), d_a.end(), d_b.begin(), d_b.end(), cuda::make_discard_iterator());
 
-  thrust::discard_iterator<> reference(h_reference.size());
+  cuda::discard_iterator reference(h_reference.size());
 
   ASSERT_EQUAL_QUIET(reference, h_result);
   ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/tabulate.cu b/thrust/testing/tabulate.cu
index e72349ffe80..82c8d0a6cb1 100644
--- a/thrust/testing/tabulate.cu
+++ b/thrust/testing/tabulate.cu
@@ -1,8 +1,9 @@
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/tabulate.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename ForwardIterator, typename UnaryOperation>
@@ -85,9 +86,7 @@ DECLARE_VARIABLE_UNITTEST(TestTabulate);
 template <typename T>
 void TestTabulateToDiscardIterator(size_t n)
 {
-  thrust::tabulate(thrust::discard_iterator<thrust::device_system_tag>(),
-                   thrust::discard_iterator<thrust::device_system_tag>(n),
-                   ::cuda::std::identity{});
+  thrust::tabulate(thrust::device, cuda::discard_iterator(), cuda::discard_iterator(n), ::cuda::std::identity{});
 
   // nothing to check -- just make sure it compiles
 }
diff --git a/thrust/testing/transform_output_iterator_reduce_by_key.cu b/thrust/testing/transform_output_iterator_reduce_by_key.cu
index 4f5b5a64e02..951698560c0 100644
--- a/thrust/testing/transform_output_iterator_reduce_by_key.cu
+++ b/thrust/testing/transform_output_iterator_reduce_by_key.cu
@@ -3,12 +3,13 @@
 #include <thrust/functional.h>
 #include <thrust/host_vector.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/transform_output_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/sequence.h>
 #include <thrust/sort.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename T>
@@ -32,7 +33,7 @@ struct TestTransformOutputIteratorReduceByKey
       h_keys.begin(),
       h_keys.end(),
       thrust::make_transform_iterator(h_values.begin(), ::cuda::std::negate<T>()),
-      thrust::discard_iterator<T>{},
+      cuda::discard_iterator{},
       h_result.begin());
     // run on device
     thrust::reduce_by_key(
@@ -40,7 +41,7 @@ struct TestTransformOutputIteratorReduceByKey
       d_keys.begin(),
       d_keys.end(),
       d_values.begin(),
-      thrust::discard_iterator<T>{},
+      cuda::discard_iterator{},
       thrust::make_transform_output_iterator(d_result.begin(), ::cuda::std::negate<T>()));
 
     ASSERT_EQUAL(h_result, d_result);
diff --git a/thrust/testing/transform_scan.cu b/thrust/testing/transform_scan.cu
index 0d0d1c56be8..0f7f14251e4 100644
--- a/thrust/testing/transform_scan.cu
+++ b/thrust/testing/transform_scan.cu
@@ -1,11 +1,12 @@
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/tabulate.h>
 #include <thrust/transform_scan.h>
 
+#include <cuda/iterator>
+
 #include <algorithm>
 #include <numeric>
 
@@ -325,20 +326,20 @@ struct TestTransformScanToDiscardIterator
     thrust::host_vector<T> h_input   = unittest::random_integers<T>(n);
     thrust::device_vector<T> d_input = h_input;
 
-    thrust::discard_iterator<> reference(n);
+    cuda::discard_iterator reference(n);
 
-    thrust::discard_iterator<> h_result = thrust::transform_inclusive_scan(
-      h_input.begin(), h_input.end(), thrust::make_discard_iterator(), ::cuda::std::negate<T>(), ::cuda::std::plus<T>());
+    cuda::discard_iterator h_result = thrust::transform_inclusive_scan(
+      h_input.begin(), h_input.end(), cuda::make_discard_iterator(), ::cuda::std::negate<T>(), ::cuda::std::plus<T>());
 
-    thrust::discard_iterator<> d_result = thrust::transform_inclusive_scan(
-      d_input.begin(), d_input.end(), thrust::make_discard_iterator(), ::cuda::std::negate<T>(), ::cuda::std::plus<T>());
+    cuda::discard_iterator d_result = thrust::transform_inclusive_scan(
+      d_input.begin(), d_input.end(), cuda::make_discard_iterator(), ::cuda::std::negate<T>(), ::cuda::std::plus<T>());
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
 
     h_result = thrust::transform_inclusive_scan(
       h_input.begin(),
       h_input.end(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       ::cuda::std::negate<T>(),
       (T) 11,
       ::cuda::std::plus<T>());
@@ -346,7 +347,7 @@ struct TestTransformScanToDiscardIterator
     d_result = thrust::transform_inclusive_scan(
       d_input.begin(),
       d_input.end(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       ::cuda::std::negate<T>(),
       (T) 11,
       ::cuda::std::plus<T>());
@@ -354,7 +355,7 @@ struct TestTransformScanToDiscardIterator
     h_result = thrust::transform_exclusive_scan(
       h_input.begin(),
       h_input.end(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       ::cuda::std::negate<T>(),
       (T) 11,
       ::cuda::std::plus<T>());
@@ -362,7 +363,7 @@ struct TestTransformScanToDiscardIterator
     d_result = thrust::transform_exclusive_scan(
       d_input.begin(),
       d_input.end(),
-      thrust::make_discard_iterator(),
+      cuda::make_discard_iterator(),
       ::cuda::std::negate<T>(),
       (T) 11,
       ::cuda::std::plus<T>());
diff --git a/thrust/testing/unique.cu b/thrust/testing/unique.cu
index fa7e41c308c..19e7cb45b06 100644
--- a/thrust/testing/unique.cu
+++ b/thrust/testing/unique.cu
@@ -1,8 +1,8 @@
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/unique.h>
 
+#include <cuda/iterator>
 #include <cuda/std/array>
 
 #include <unittest/unittest.h>
@@ -238,15 +238,13 @@ struct TestUniqueCopyToDiscardIterator
     thrust::host_vector<T> h_unique = h_data;
     h_unique.erase(thrust::unique(h_unique.begin(), h_unique.end()), h_unique.end());
 
-    thrust::discard_iterator<> reference(h_unique.size());
+    cuda::discard_iterator reference(h_unique.size());
 
     typename thrust::device_vector<T>::iterator d_new_last;
 
-    thrust::discard_iterator<> h_result =
-      thrust::unique_copy(h_data.begin(), h_data.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator h_result = thrust::unique_copy(h_data.begin(), h_data.end(), cuda::make_discard_iterator());
 
-    thrust::discard_iterator<> d_result =
-      thrust::unique_copy(d_data.begin(), d_data.end(), thrust::make_discard_iterator());
+    cuda::discard_iterator d_result = thrust::unique_copy(d_data.begin(), d_data.end(), cuda::make_discard_iterator());
 
     ASSERT_EQUAL_QUIET(reference, h_result);
     ASSERT_EQUAL_QUIET(reference, d_result);
diff --git a/thrust/testing/unique_by_key.cu b/thrust/testing/unique_by_key.cu
index 24de2ed2bda..51e2589107b 100644
--- a/thrust/testing/unique_by_key.cu
+++ b/thrust/testing/unique_by_key.cu
@@ -1,8 +1,9 @@
 #include <thrust/functional.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/retag.h>
 #include <thrust/unique.h>
 
+#include <cuda/iterator>
+
 #include <unittest/unittest.h>
 
 template <typename ValueT>
@@ -335,51 +336,51 @@ struct TestUniqueCopyByKeyToDiscardIterator
     size_t num_unique_keys = h_unique_keys.size();
 
     // mask both outputs
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> h_result1 = thrust::unique_by_key_copy(
-      h_keys.begin(), h_keys.end(), h_vals.begin(), thrust::make_discard_iterator(), thrust::make_discard_iterator());
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> h_result1 = thrust::unique_by_key_copy(
+      h_keys.begin(), h_keys.end(), h_vals.begin(), cuda::make_discard_iterator(), cuda::make_discard_iterator());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> d_result1 = thrust::unique_by_key_copy(
-      d_keys.begin(), d_keys.end(), d_vals.begin(), thrust::make_discard_iterator(), thrust::make_discard_iterator());
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> d_result1 = thrust::unique_by_key_copy(
+      d_keys.begin(), d_keys.end(), d_vals.begin(), cuda::make_discard_iterator(), cuda::make_discard_iterator());
 
-    cuda::std::pair<thrust::discard_iterator<>, thrust::discard_iterator<>> reference1 = cuda::std::make_pair(
-      thrust::make_discard_iterator(num_unique_keys), thrust::make_discard_iterator(num_unique_keys));
+    cuda::std::pair<cuda::discard_iterator, cuda::discard_iterator> reference1 =
+      cuda::std::make_pair(cuda::make_discard_iterator(num_unique_keys), cuda::make_discard_iterator(num_unique_keys));
 
     ASSERT_EQUAL_QUIET(reference1, h_result1);
     ASSERT_EQUAL_QUIET(reference1, d_result1);
 
     // mask values output
-    cuda::std::pair<typename thrust::host_vector<K>::iterator, thrust::discard_iterator<>> h_result2 =
+    cuda::std::pair<typename thrust::host_vector<K>::iterator, cuda::discard_iterator> h_result2 =
       thrust::unique_by_key_copy(
-        h_keys.begin(), h_keys.end(), h_vals.begin(), h_keys_output.begin(), thrust::make_discard_iterator());
+        h_keys.begin(), h_keys.end(), h_vals.begin(), h_keys_output.begin(), cuda::make_discard_iterator());
 
-    cuda::std::pair<typename thrust::device_vector<K>::iterator, thrust::discard_iterator<>> d_result2 =
+    cuda::std::pair<typename thrust::device_vector<K>::iterator, cuda::discard_iterator> d_result2 =
       thrust::unique_by_key_copy(
-        d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys_output.begin(), thrust::make_discard_iterator());
+        d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys_output.begin(), cuda::make_discard_iterator());
 
-    cuda::std::pair<typename thrust::host_vector<K>::iterator, thrust::discard_iterator<>> h_reference2 =
-      cuda::std::make_pair(h_keys_output.begin() + num_unique_keys, thrust::make_discard_iterator(num_unique_keys));
+    cuda::std::pair<typename thrust::host_vector<K>::iterator, cuda::discard_iterator> h_reference2 =
+      cuda::std::make_pair(h_keys_output.begin() + num_unique_keys, cuda::make_discard_iterator(num_unique_keys));
 
-    cuda::std::pair<typename thrust::device_vector<K>::iterator, thrust::discard_iterator<>> d_reference2 =
-      cuda::std::make_pair(d_keys_output.begin() + num_unique_keys, thrust::make_discard_iterator(num_unique_keys));
+    cuda::std::pair<typename thrust::device_vector<K>::iterator, cuda::discard_iterator> d_reference2 =
+      cuda::std::make_pair(d_keys_output.begin() + num_unique_keys, cuda::make_discard_iterator(num_unique_keys));
 
     ASSERT_EQUAL(h_keys_output, d_keys_output);
     ASSERT_EQUAL_QUIET(h_reference2, h_result2);
     ASSERT_EQUAL_QUIET(d_reference2, d_result2);
 
     // mask keys output
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<V>::iterator> h_result3 =
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<V>::iterator> h_result3 =
       thrust::unique_by_key_copy(
-        h_keys.begin(), h_keys.end(), h_vals.begin(), thrust::make_discard_iterator(), h_vals_output.begin());
+        h_keys.begin(), h_keys.end(), h_vals.begin(), cuda::make_discard_iterator(), h_vals_output.begin());
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<V>::iterator> d_result3 =
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<V>::iterator> d_result3 =
       thrust::unique_by_key_copy(
-        d_keys.begin(), d_keys.end(), d_vals.begin(), thrust::make_discard_iterator(), d_vals_output.begin());
+        d_keys.begin(), d_keys.end(), d_vals.begin(), cuda::make_discard_iterator(), d_vals_output.begin());
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::host_vector<V>::iterator> h_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(num_unique_keys), h_vals_output.begin() + num_unique_keys);
+    cuda::std::pair<cuda::discard_iterator, typename thrust::host_vector<V>::iterator> h_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(num_unique_keys), h_vals_output.begin() + num_unique_keys);
 
-    cuda::std::pair<thrust::discard_iterator<>, typename thrust::device_vector<V>::iterator> d_reference3 =
-      cuda::std::make_pair(thrust::make_discard_iterator(num_unique_keys), d_vals_output.begin() + num_unique_keys);
+    cuda::std::pair<cuda::discard_iterator, typename thrust::device_vector<V>::iterator> d_reference3 =
+      cuda::std::make_pair(cuda::make_discard_iterator(num_unique_keys), d_vals_output.begin() + num_unique_keys);
 
     ASSERT_EQUAL(h_vals_output, d_vals_output);
     ASSERT_EQUAL_QUIET(h_reference3, h_result3);
@@ -437,8 +438,8 @@ struct TestUniqueCopyByKeyLargeOutCount
     auto values_in = thrust::make_counting_iterator(0ULL);
 
     // Run test
-    auto keys_out   = thrust::make_discard_iterator();
-    auto values_out = thrust::make_discard_iterator();
+    auto keys_out   = cuda::make_discard_iterator();
+    auto values_out = cuda::make_discard_iterator();
     const auto selected_aut_end =
       thrust::unique_by_key_copy(thrust::device, keys_in, keys_in + num_items, values_in, keys_out, values_out);
 
diff --git a/thrust/thrust/detail/internal_functional.h b/thrust/thrust/detail/internal_functional.h
index 5782caea033..a76a3c1f1b3 100644
--- a/thrust/thrust/detail/internal_functional.h
+++ b/thrust/thrust/detail/internal_functional.h
@@ -78,13 +78,13 @@ struct new_value_if_f
   template <class T>
   _CCCL_DEVICE_API OutputType operator()(T const& x)
   {
-    return pred(x) ? new_value : x;
+    return OutputType{pred(x) ? new_value : x};
   }
 
   template <class T, class P>
   _CCCL_DEVICE_API OutputType operator()(T const& x, P const& y)
   {
-    return pred(y) ? new_value : x;
+    return OutputType{pred(y) ? new_value : x};
   }
 };
 
diff --git a/thrust/thrust/iterator/discard_iterator.h b/thrust/thrust/iterator/discard_iterator.h
index cd4e4bb5157..215ba1a3111 100644
--- a/thrust/thrust/iterator/discard_iterator.h
+++ b/thrust/thrust/iterator/discard_iterator.h
@@ -98,8 +98,10 @@ _CCCL_DIAG_SUPPRESS_MSVC(4244 4267) // possible loss of data
  *     .. versionadded:: 2.2.0
  *  \endverbatim
  */
+//! \deprecated since 3.4
 template <typename System = use_default>
-class discard_iterator : public detail::make_discard_iterator_base<System>::type
+class CCCL_DEPRECATED_BECAUSE("Use cuda::discard_iterator instead") discard_iterator
+    : public detail::make_discard_iterator_base<System>::type
 {
   //! \cond
   friend class iterator_core_access;
@@ -140,8 +142,10 @@ class discard_iterator : public detail::make_discard_iterator_base<System>::type
 //! parameter is \c 0.
 //! \return A new \p discard_iterator with index as given by \p i.
 //! \see constant_iterator
+//! \deprecated since 3.4
 template <typename System = use_default>
-inline _CCCL_HOST_DEVICE discard_iterator<System> make_discard_iterator(::cuda::std::ptrdiff_t i = 0)
+inline CCCL_DEPRECATED_BECAUSE("Use cuda::make_discard_iterator instead") _CCCL_HOST_DEVICE discard_iterator<System>
+make_discard_iterator(::cuda::std::ptrdiff_t i = 0)
 {
   return discard_iterator<System>{i};
 }
diff --git a/thrust/thrust/iterator/permutation_iterator.h b/thrust/thrust/iterator/permutation_iterator.h
index ce9160f11ee..97fbcdbfbb5 100644
--- a/thrust/thrust/iterator/permutation_iterator.h
+++ b/thrust/thrust/iterator/permutation_iterator.h
@@ -32,6 +32,9 @@
 #include <thrust/iterator/iterator_adaptor.h>
 #include <thrust/iterator/iterator_traits.h>
 
+#include <cuda/std/__type_traits/is_same.h>
+#include <cuda/std/__type_traits/remove_cvref.h>
+
 THRUST_NAMESPACE_BEGIN
 
 template <typename, typename>
@@ -45,13 +48,24 @@ struct make_permutation_iterator_base
   using System1 = iterator_system_t<ElementIterator>;
   using System2 = iterator_system_t<IndexIterator>;
 
+  // We need this to make proxy iterators work because those have a void reference/value type
+  using iterator_value_t =
+    ::cuda::std::conditional_t<::cuda::std::is_same_v<it_value_t<ElementIterator>, void>,
+                               ::cuda::std::remove_cvref_t<decltype(*::cuda::std::declval<ElementIterator>())>,
+                               it_value_t<ElementIterator>>;
+
+  using iterator_reference_t =
+    ::cuda::std::conditional_t<::cuda::std::is_same_v<it_reference_t<ElementIterator>, void>,
+                               decltype(*::cuda::std::declval<ElementIterator>()),
+                               it_reference_t<ElementIterator>>;
+
   using type =
     iterator_adaptor<permutation_iterator<ElementIterator, IndexIterator>,
                      IndexIterator,
-                     it_value_t<ElementIterator>,
+                     iterator_value_t,
                      minimum_system_t<System1, System2>,
                      use_default,
-                     it_reference_t<ElementIterator>>;
+                     iterator_reference_t>;
 };
 } // namespace detail
 
diff --git a/thrust/thrust/iterator/zip_iterator.h b/thrust/thrust/iterator/zip_iterator.h
index 11dee61b6a7..89e9390a1ce 100644
--- a/thrust/thrust/iterator/zip_iterator.h
+++ b/thrust/thrust/iterator/zip_iterator.h
@@ -34,10 +34,12 @@
 
 #include <cuda/std/__iterator/advance.h>
 #include <cuda/std/__iterator/distance.h>
+#include <cuda/std/__iterator/readable_traits.h>
 #include <cuda/std/__type_traits/conditional.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_constructible.h>
-#include <cuda/std/__type_traits/is_same.h>
+#include <cuda/std/__type_traits/is_void.h>
+#include <cuda/std/__type_traits/remove_cvref.h>
 #include <cuda/std/__utility/declval.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/tuple>
@@ -61,15 +63,24 @@ struct make_zip_iterator_base<::cuda::std::tuple<Its...>>
   // We need this to make proxy iterators work because those have a void reference type
   template <class Iter>
   using zip_iterator_reference_t =
-    ::cuda::std::conditional_t<::cuda::std::is_same_v<it_reference_t<Iter>, void>,
+    ::cuda::std::conditional_t<::cuda::std::is_void_v<it_reference_t<Iter>>,
                                decltype(*::cuda::std::declval<Iter>()),
                                it_reference_t<Iter>>;
 
   // reference type is the type of the tuple obtained from the iterator's reference types.
   using reference = tuple_of_iterator_references<zip_iterator_reference_t<Its>...>;
 
+  // We need this to make proxy iterators work because those have a void value type. We also
+  // cannot use cuda::std::iter_value_t directly because it does not consider nested
+  // value_type's.
+  template <class Iter>
+  using zip_iterator_value_t =
+    ::cuda::std::_If<::cuda::std::is_void_v<it_value_t<Iter>>,
+                     ::cuda::std::remove_cvref_t<decltype(*::cuda::std::declval<Iter>())>,
+                     it_value_t<Iter>>;
+
   // Boost's Value type is the same as reference type. using value_type = reference;
-  using value_type = ::cuda::std::tuple<it_value_t<Its>...>;
+  using value_type = ::cuda::std::tuple<zip_iterator_value_t<Its>...>;
 
   // Difference type is the first iterator's difference type
   using difference_type = it_difference_t<::cuda::std::tuple_element_t<0, ::cuda::std::tuple<Its...>>>;
diff --git a/thrust/thrust/system/cuda/detail/replace.h b/thrust/thrust/system/cuda/detail/replace.h
index 79160764729..fec67ab1170 100644
--- a/thrust/thrust/system/cuda/detail/replace.h
+++ b/thrust/thrust/system/cuda/detail/replace.h
@@ -64,7 +64,7 @@ OutputIt _CCCL_HOST_DEVICE replace_copy_if(
   Predicate predicate,
   T const& new_value)
 {
-  using output_type    = thrust::detail::it_value_t<OutputIt>;
+  using output_type    = ::cuda::std::iter_value_t<OutputIt>;
   using new_value_if_t = thrust::detail::new_value_if_f<Predicate, T, output_type>;
   return cuda_cub::transform(policy, first, last, result, new_value_if_t{predicate, new_value});
 }
@@ -79,7 +79,7 @@ OutputIt _CCCL_HOST_DEVICE replace_copy_if(
   Predicate predicate,
   T const& new_value)
 {
-  using output_type    = thrust::detail::it_value_t<OutputIt>;
+  using output_type    = ::cuda::std::iter_value_t<OutputIt>;
   using new_value_if_t = thrust::detail::new_value_if_f<Predicate, T, output_type>;
   return cuda_cub::transform(policy, first, last, stencil, result, new_value_if_t{predicate, new_value});
 }
diff --git a/thrust/thrust/system/detail/generic/replace.inl b/thrust/thrust/system/detail/generic/replace.inl
index fb13bef2459..c66d13068e3 100644
--- a/thrust/thrust/system/detail/generic/replace.inl
+++ b/thrust/thrust/system/detail/generic/replace.inl
@@ -17,6 +17,8 @@
 #include <thrust/system/detail/generic/replace.h>
 #include <thrust/transform.h>
 
+#include <cuda/std/__iterator/readable_traits.h>
+
 THRUST_NAMESPACE_BEGIN
 namespace system::detail::generic
 {
@@ -35,7 +37,7 @@ struct new_value_if
   template <typename InputType>
   _CCCL_HOST_DEVICE OutputType operator()(const InputType& x) const
   {
-    return pred(x) ? new_value : x;
+    return OutputType{pred(x) ? new_value : x};
   } // end operator()()
 
   // this version of operator()() works like the previous but
@@ -43,7 +45,7 @@ struct new_value_if
   template <typename InputType, typename PredicateArgumentType>
   _CCCL_HOST_DEVICE OutputType operator()(const InputType& x, const PredicateArgumentType& y)
   {
-    return pred(y) ? new_value : x;
+    return OutputType{pred(y) ? new_value : x};
   } // end operator()()
 
   Predicate pred;
@@ -77,7 +79,7 @@ _CCCL_HOST_DEVICE OutputIterator replace_copy_if(
   Predicate pred,
   const T& new_value)
 {
-  using OutputType = thrust::detail::it_value_t<OutputIterator>;
+  using OutputType = ::cuda::std::iter_value_t<OutputIterator>;
 
   detail::new_value_if<Predicate, T, OutputType> op(pred, new_value);
   return thrust::transform(exec, first, last, result, op);
@@ -98,7 +100,7 @@ _CCCL_HOST_DEVICE OutputIterator replace_copy_if(
   Predicate pred,
   const T& new_value)
 {
-  using OutputType = thrust::detail::it_value_t<OutputIterator>;
+  using OutputType = ::cuda::std::iter_value_t<OutputIterator>;
 
   detail::new_value_if<Predicate, T, OutputType> op(pred, new_value);
   return thrust::transform(exec, first, last, stencil, result, op);
diff --git a/thrust/thrust/system/detail/generic/shuffle.inl b/thrust/thrust/system/detail/generic/shuffle.inl
index 50cdf5acffb..1d6a7a09e22 100644
--- a/thrust/thrust/system/detail/generic/shuffle.inl
+++ b/thrust/thrust/system/detail/generic/shuffle.inl
@@ -5,13 +5,14 @@
 
 #include <thrust/detail/random_bijection.h>
 #include <thrust/detail/temporary_array.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/iterator/transform_output_iterator.h>
 #include <thrust/random.h>
 #include <thrust/scan.h>
 #include <thrust/system/detail/generic/shuffle.h>
 
+#include <cuda/__iterator/counting_iterator.h>
+#include <cuda/__iterator/discard_iterator.h>
 #include <cuda/std/cstdint>
 
 THRUST_NAMESPACE_BEGIN
@@ -94,12 +95,11 @@ _CCCL_HOST_DEVICE void shuffle_copy(
 
   // perform stream compaction over length n bijection to get length m
   // pseudorandom bijection over the original input
-  thrust::counting_iterator<std::uint64_t> indices(0);
+  ::cuda::counting_iterator<std::uint64_t> indices(0);
   thrust::transform_iterator<construct_key_flag_op, decltype(indices), key_flag_tuple> key_flag_it(
     indices, construct_key_flag_op(m, bijection));
   write_output_op<RandomIterator, decltype(result)> write_functor{m, first, result};
-  auto gather_output_it =
-    thrust::make_transform_output_iterator(thrust::discard_iterator<std::size_t>(), write_functor);
+  auto gather_output_it = thrust::make_transform_output_iterator(::cuda::discard_iterator(), write_functor);
   // the feistel_bijection outputs a stream of permuted indices in range [0,n)
   // flag each value < m and compact it, so we have a set of permuted indices in
   // range [0,m) each thread gathers an input element according to its
diff --git a/thrust/thrust/system/detail/sequential/unique_by_key.h b/thrust/thrust/system/detail/sequential/unique_by_key.h
index db1cb87f470..6e5b13cfd6e 100644
--- a/thrust/thrust/system/detail/sequential/unique_by_key.h
+++ b/thrust/thrust/system/detail/sequential/unique_by_key.h
@@ -40,18 +40,23 @@ _CCCL_HOST_DEVICE ::cuda::std::pair<OutputIterator1, OutputIterator2> unique_by_
   OutputIterator2 values_output,
   BinaryPredicate binary_pred)
 {
-  using InputKeyType    = thrust::detail::it_value_t<InputIterator1>;
-  using OutputValueType = thrust::detail::it_value_t<OutputIterator2>;
+  using InputKeyType = ::cuda::std::iter_value_t<InputIterator1>;
 
   if (keys_first != keys_last)
   {
-    InputKeyType temp_key      = *keys_first;
-    OutputValueType temp_value = *values_first;
+    InputKeyType temp_key = *keys_first;
+    // Store by value, and do not convert to OutputIterator::value_type via:
+    //
+    // OutputIteratorValueType temp_value = *values_first;
+    //
+    // Iterators only guarantee that assignment between iterator value-types is valid, not
+    // construction.
+    auto temp_value = *values_first;
 
     for (++keys_first, ++values_first; keys_first != keys_last; ++keys_first, (void) ++values_first)
     {
-      InputKeyType key      = *keys_first;
-      OutputValueType value = *values_first;
+      InputKeyType key = *keys_first;
+      auto value       = *values_first;
 
       if (!binary_pred(temp_key, key))
       {