From 7e427c3c9512d8e0811ac99df5b4ea3172e1ae01 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Tue, 28 Jan 2025 12:20:11 +0100
Subject: [PATCH 1/2] Deprecate and replace some thrust iterator traits

Deprecate and remove use of thrust::
* iterator_traits
* iterator_value
* iterator_value_t
* iterator_pointer
* iterator_pointer_t
* iterator_reference
* iterator_reference_t
* iterator_difference
* iterator_difference_t

Add as replacement thrust::detail::
* it_value_t
* it_reference_t
* it_difference_t
* it_pointer_t
---
 thrust/examples/cuda/range_view.cu            | 10 ++-
 thrust/examples/expand.cu                     |  2 +-
 thrust/examples/repeated_range.cu             |  2 +-
 thrust/examples/strided_range.cu              |  2 +-
 thrust/examples/tiled_range.cu                |  2 +-
 thrust/testing/counting_iterator.cu           |  2 +-
 thrust/testing/cuda/adjacent_difference.cu    |  4 +-
 thrust/testing/omp/reduce_intervals.cu        |  2 +-
 thrust/testing/sort_permutation_iterator.cu   |  2 +-
 thrust/testing/trivial_sequence.cu            |  2 +-
 thrust/testing/unique.cu                      |  5 +-
 thrust/testing/unittest/assertions.h          |  8 +-
 thrust/testing/zip_iterator.cu                |  4 +-
 thrust/thrust/advance.h                       |  4 +-
 thrust/thrust/count.h                         |  9 +--
 thrust/thrust/detail/advance.inl              |  7 +-
 .../detail/allocator/copy_construct_range.inl | 10 +--
 .../detail/allocator/tagged_allocator.h       |  4 +-
 thrust/thrust/detail/count.h                  |  9 +--
 thrust/thrust/detail/count.inl                |  9 +--
 thrust/thrust/detail/distance.inl             |  3 +-
 thrust/thrust/detail/get_iterator_value.h     |  3 +-
 thrust/thrust/detail/memory_algorithms.h      | 12 +--
 thrust/thrust/detail/overlapped_copy.h        |  2 +-
 thrust/thrust/detail/pointer.h                | 22 +-----
 thrust/thrust/detail/range/head_flags.h       |  4 +-
 thrust/thrust/detail/range/tail_flags.h       |  4 +-
 thrust/thrust/detail/reduce.inl               |  4 +-
 thrust/thrust/detail/temporary_array.h        |  2 +-
 thrust/thrust/detail/trivial_sequence.h       |  2 +-
 .../type_traits/iterator/is_output_iterator.h | 30 ++------
 thrust/thrust/detail/unique.inl               |  8 +-
 thrust/thrust/detail/vector_base.inl          |  2 +-
 thrust/thrust/distance.h                      |  2 +-
 thrust/thrust/iterator/counting_iterator.h    |  2 +-
 .../iterator/detail/iterator_adaptor_base.h   |  6 +-
 .../thrust/iterator/detail/tagged_iterator.h  |  6 +-
 thrust/thrust/iterator/iterator_traits.h      | 77 ++++++++++++++-----
 thrust/thrust/iterator/permutation_iterator.h |  4 +-
 .../transform_input_output_iterator.h         |  4 +-
 thrust/thrust/iterator/transform_iterator.h   | 10 +--
 thrust/thrust/iterator/zip_iterator.h         | 25 +++++-
 thrust/thrust/reduce.h                        |  4 +-
 .../system/cuda/detail/adjacent_difference.h  |  6 +-
 .../system/cuda/detail/async/inclusive_scan.h |  6 +-
 thrust/thrust/system/cuda/detail/copy_if.h    |  2 +-
 .../system/cuda/detail/core/load_iterator.h   |  4 +-
 thrust/thrust/system/cuda/detail/core/util.h  |  2 +-
 thrust/thrust/system/cuda/detail/count.h      |  6 +-
 thrust/thrust/system/cuda/detail/equal.h      |  2 +-
 thrust/thrust/system/cuda/detail/extrema.h    | 15 ++--
 thrust/thrust/system/cuda/detail/find.h       |  6 +-
 thrust/thrust/system/cuda/detail/for_each.h   |  2 +-
 thrust/thrust/system/cuda/detail/get_value.h  |  4 +-
 .../cuda/detail/internal/copy_cross_system.h  |  6 +-
 .../detail/internal/copy_device_to_device.h   |  2 +-
 thrust/thrust/system/cuda/detail/merge.h      |  4 +-
 thrust/thrust/system/cuda/detail/mismatch.h   |  4 +-
 thrust/thrust/system/cuda/detail/partition.h  |  4 +-
 thrust/thrust/system/cuda/detail/reduce.h     |  7 +-
 .../thrust/system/cuda/detail/reduce_by_key.h | 16 ++--
 thrust/thrust/system/cuda/detail/replace.h    |  4 +-
 thrust/thrust/system/cuda/detail/reverse.h    |  4 +-
 thrust/thrust/system/cuda/detail/scan.h       | 10 +--
 .../thrust/system/cuda/detail/scan_by_key.h   |  4 +-
 .../system/cuda/detail/set_operations.h       | 38 ++++-----
 thrust/thrust/system/cuda/detail/sort.h       | 18 ++---
 .../thrust/system/cuda/detail/swap_ranges.h   |  6 +-
 thrust/thrust/system/cuda/detail/tabulate.h   |  2 +-
 thrust/thrust/system/cuda/detail/transform.h  |  8 +-
 .../system/cuda/detail/transform_reduce.h     |  3 +-
 .../system/cuda/detail/transform_scan.h       | 10 +--
 .../system/cuda/detail/uninitialized_copy.h   |  4 +-
 .../system/cuda/detail/uninitialized_fill.h   |  2 +-
 thrust/thrust/system/cuda/detail/unique.h     | 12 +--
 .../thrust/system/cuda/detail/unique_by_key.h |  6 +-
 .../detail/generic/adjacent_difference.inl    |  4 +-
 .../system/detail/generic/binary_search.inl   |  8 +-
 .../thrust/system/detail/generic/copy_if.inl  |  2 +-
 thrust/thrust/system/detail/generic/count.h   |  4 +-
 thrust/thrust/system/detail/generic/count.inl |  8 +-
 .../thrust/system/detail/generic/distance.h   |  2 +-
 .../thrust/system/detail/generic/distance.inl |  9 +--
 .../thrust/system/detail/generic/extrema.inl  | 18 ++---
 thrust/thrust/system/detail/generic/find.inl  |  2 +-
 .../thrust/system/detail/generic/generate.inl |  6 +-
 thrust/thrust/system/detail/generic/merge.inl |  4 +-
 .../system/detail/generic/partition.inl       |  9 +--
 thrust/thrust/system/detail/generic/reduce.h  |  2 +-
 .../thrust/system/detail/generic/reduce.inl   |  4 +-
 .../system/detail/generic/reduce_by_key.inl   | 16 ++--
 .../thrust/system/detail/generic/remove.inl   |  4 +-
 .../thrust/system/detail/generic/replace.inl  |  4 +-
 .../thrust/system/detail/generic/reverse.inl  |  2 +-
 .../detail/generic/scalar/binary_search.inl   |  4 +-
 thrust/thrust/system/detail/generic/scan.inl  |  3 +-
 .../system/detail/generic/scan_by_key.inl     |  4 +-
 .../thrust/system/detail/generic/sequence.inl |  2 +-
 .../system/detail/generic/set_operations.inl  | 18 ++---
 .../thrust/system/detail/generic/shuffle.inl  |  2 +-
 thrust/thrust/system/detail/generic/sort.inl  | 10 +--
 .../thrust/system/detail/generic/tabulate.inl |  2 +-
 .../system/detail/generic/transform_scan.inl  |  5 +-
 .../detail/generic/uninitialized_copy.inl     | 14 ++--
 .../detail/generic/uninitialized_fill.inl     |  8 +-
 thrust/thrust/system/detail/generic/unique.h  |  4 +-
 .../thrust/system/detail/generic/unique.inl   | 12 +--
 .../system/detail/generic/unique_by_key.inl   | 10 +--
 .../detail/sequential/adjacent_difference.h   |  2 +-
 .../system/detail/sequential/binary_search.h  |  4 +-
 .../thrust/system/detail/sequential/copy.inl  |  2 +-
 .../system/detail/sequential/general_copy.h   | 28 ++++---
 .../system/detail/sequential/get_value.h      |  2 +-
 .../system/detail/sequential/insertion_sort.h |  6 +-
 .../system/detail/sequential/partition.h      |  6 +-
 .../system/detail/sequential/reduce_by_key.h  |  6 +-
 thrust/thrust/system/detail/sequential/scan.h |  7 +-
 .../system/detail/sequential/scan_by_key.h    |  6 +-
 .../thrust/system/detail/sequential/sort.inl  |  8 +-
 .../detail/sequential/stable_merge_sort.inl   | 17 ++--
 .../sequential/stable_primitive_sort.inl      |  4 +-
 .../detail/sequential/stable_radix_sort.inl   | 16 ++--
 .../thrust/system/detail/sequential/unique.h  |  8 +-
 .../system/detail/sequential/unique_by_key.h  |  4 +-
 thrust/thrust/system/omp/detail/for_each.inl  |  2 +-
 thrust/thrust/system/omp/detail/reduce.inl    |  2 +-
 .../system/omp/detail/reduce_intervals.inl    |  2 +-
 thrust/thrust/system/omp/detail/sort.inl      | 10 +--
 thrust/thrust/system/omp/detail/unique.h      |  2 +-
 thrust/thrust/system/omp/detail/unique.inl    |  2 +-
 thrust/thrust/system/tbb/detail/copy_if.inl   |  2 +-
 thrust/thrust/system/tbb/detail/reduce.inl    |  2 +-
 .../system/tbb/detail/reduce_by_key.inl       | 22 +++---
 .../system/tbb/detail/reduce_intervals.h      |  4 +-
 thrust/thrust/system/tbb/detail/scan.inl      | 13 ++--
 thrust/thrust/system/tbb/detail/sort.inl      | 10 +--
 thrust/thrust/system/tbb/detail/unique.h      |  2 +-
 thrust/thrust/system/tbb/detail/unique.inl    |  2 +-
 .../type_traits/is_trivially_relocatable.h    |  9 +--
 thrust/thrust/unique.h                        |  8 +-
 140 files changed, 511 insertions(+), 481 deletions(-)

diff --git a/thrust/examples/cuda/range_view.cu b/thrust/examples/cuda/range_view.cu
index f045d67bca7..5b7e3825d79 100644
--- a/thrust/examples/cuda/range_view.cu
+++ b/thrust/examples/cuda/range_view.cu
@@ -3,6 +3,8 @@
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
 
+#include <cuda/std/iterator>
+
 #include <iostream>
 
 #include "../include/host_device.h"
@@ -20,10 +22,10 @@ class range_view
 {
 public:
   using iterator        = Iterator;
-  using value_type      = typename thrust::iterator_traits<iterator>::value_type;
-  using pointer         = typename thrust::iterator_traits<iterator>::pointer;
-  using difference_type = typename thrust::iterator_traits<iterator>::difference_type;
-  using reference       = typename thrust::iterator_traits<iterator>::reference;
+  using value_type      = typename cuda::std::iterator_traits<iterator>::value_type;
+  using pointer         = typename cuda::std::iterator_traits<iterator>::pointer;
+  using difference_type = typename cuda::std::iterator_traits<iterator>::difference_type;
+  using reference       = typename cuda::std::iterator_traits<iterator>::reference;
 
 private:
   const iterator first;
diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu
index ecded38ffd8..350ec1d4099 100644
--- a/thrust/examples/expand.cu
+++ b/thrust/examples/expand.cu
@@ -20,7 +20,7 @@
 template <typename InputIterator1, typename InputIterator2, typename OutputIterator>
 OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, OutputIterator output)
 {
-  using difference_type = typename thrust::iterator_difference<InputIterator1>::type;
+  using difference_type = thrust::detail::it_difference_t<InputIterator1>;
 
   difference_type input_size  = thrust::distance(first1, last1);
   difference_type output_size = thrust::reduce(first1, last1);
diff --git a/thrust/examples/repeated_range.cu b/thrust/examples/repeated_range.cu
index c86bb22f64e..debd13c58f9 100644
--- a/thrust/examples/repeated_range.cu
+++ b/thrust/examples/repeated_range.cu
@@ -21,7 +21,7 @@ template <typename Iterator>
 class repeated_range
 {
 public:
-  using difference_type = typename thrust::iterator_difference<Iterator>::type;
+  using difference_type = thrust::detail::it_difference_t<Iterator>;
 
   struct repeat_functor
   {
diff --git a/thrust/examples/strided_range.cu b/thrust/examples/strided_range.cu
index 40ad38fcb9a..dbdbac82d62 100644
--- a/thrust/examples/strided_range.cu
+++ b/thrust/examples/strided_range.cu
@@ -21,7 +21,7 @@ template <typename Iterator>
 class strided_range
 {
 public:
-  using difference_type = typename thrust::iterator_difference<Iterator>::type;
+  using difference_type = thrust::detail::it_difference_t<Iterator>;
 
   struct stride_functor
   {
diff --git a/thrust/examples/tiled_range.cu b/thrust/examples/tiled_range.cu
index ca69dea2f77..220591dc8f5 100644
--- a/thrust/examples/tiled_range.cu
+++ b/thrust/examples/tiled_range.cu
@@ -21,7 +21,7 @@ template <typename Iterator>
 class tiled_range
 {
 public:
-  using difference_type = typename thrust::iterator_difference<Iterator>::type;
+  using difference_type = thrust::detail::it_difference_t<Iterator>;
 
   struct tile_functor
   {
diff --git a/thrust/testing/counting_iterator.cu b/thrust/testing/counting_iterator.cu
index 69905636c43..24371d12de6 100644
--- a/thrust/testing/counting_iterator.cu
+++ b/thrust/testing/counting_iterator.cu
@@ -235,7 +235,7 @@ DECLARE_UNITTEST(TestCountingIteratorLowerBound);
 void TestCountingIteratorDifference()
 {
   using Iterator   = thrust::counting_iterator<std::uint64_t>;
-  using Difference = thrust::iterator_difference<Iterator>::type;
+  using Difference = thrust::detail::it_difference_t<Iterator>;
 
   Difference diff = std::numeric_limits<std::uint32_t>::max() + 1;
 
diff --git a/thrust/testing/cuda/adjacent_difference.cu b/thrust/testing/cuda/adjacent_difference.cu
index d05defe5a58..cbc8fcd515e 100644
--- a/thrust/testing/cuda/adjacent_difference.cu
+++ b/thrust/testing/cuda/adjacent_difference.cu
@@ -96,9 +96,9 @@ DECLARE_UNITTEST(TestAdjacentDifferenceCudaStreams);
 struct detect_wrong_difference
 {
   using difference_type   = void;
-  using value_type        = void;
+  using value_type        = long long;
   using pointer           = void;
-  using reference         = void;
+  using reference         = detect_wrong_difference;
   using iterator_category = ::cuda::std::output_iterator_tag;
 
   bool* flag;
diff --git a/thrust/testing/omp/reduce_intervals.cu b/thrust/testing/omp/reduce_intervals.cu
index 309faf69cc9..c3995570a3b 100644
--- a/thrust/testing/omp/reduce_intervals.cu
+++ b/thrust/testing/omp/reduce_intervals.cu
@@ -8,7 +8,7 @@
 template <typename InputIterator, typename OutputIterator, typename BinaryFunction, typename Decomposition>
 void reduce_intervals(InputIterator input, OutputIterator output, BinaryFunction binary_op, Decomposition decomp)
 {
-  using OutputType = typename thrust::iterator_value<OutputIterator>::type;
+  using OutputType = thrust::detail::it_value_t<OutputIterator>;
   using index_type = typename Decomposition::index_type;
 
   // wrap binary_op
diff --git a/thrust/testing/sort_permutation_iterator.cu b/thrust/testing/sort_permutation_iterator.cu
index d17281be5a0..0e51a00a58b 100644
--- a/thrust/testing/sort_permutation_iterator.cu
+++ b/thrust/testing/sort_permutation_iterator.cu
@@ -10,7 +10,7 @@ template <typename Iterator>
 class strided_range
 {
 public:
-  using difference_type = typename thrust::iterator_difference<Iterator>::type;
+  using difference_type = thrust::detail::it_difference_t<Iterator>;
 
   struct stride_functor
   {
diff --git a/thrust/testing/trivial_sequence.cu b/thrust/testing/trivial_sequence.cu
index 7438b7e406a..b9072bcb382 100644
--- a/thrust/testing/trivial_sequence.cu
+++ b/thrust/testing/trivial_sequence.cu
@@ -10,7 +10,7 @@ void test(Iterator first, Iterator last)
   using System = typename thrust::iterator_system<Iterator>::type;
   System system;
   thrust::detail::trivial_sequence<Iterator, System> ts(system, first, last);
-  using ValueType = typename thrust::iterator_traits<Iterator>::value_type;
+  using ValueType = typename ::cuda::std::iterator_traits<Iterator>::value_type;
 
   ASSERT_EQUAL_QUIET((ValueType) ts.begin()[0], ValueType(0, 11));
   ASSERT_EQUAL_QUIET((ValueType) ts.begin()[1], ValueType(2, 11));
diff --git a/thrust/testing/unique.cu b/thrust/testing/unique.cu
index 8a03353b517..fa7e41c308c 100644
--- a/thrust/testing/unique.cu
+++ b/thrust/testing/unique.cu
@@ -79,7 +79,7 @@ void TestUniqueCopyDispatchImplicit()
 DECLARE_UNITTEST(TestUniqueCopyDispatchImplicit);
 
 template <typename ForwardIterator>
-typename thrust::iterator_traits<ForwardIterator>::difference_type
+typename ::cuda::std::iterator_traits<ForwardIterator>::difference_type
 unique_count(my_system& system, ForwardIterator, ForwardIterator)
 {
   system.validate_dispatch();
@@ -98,7 +98,8 @@ void TestUniqueCountDispatchExplicit()
 DECLARE_UNITTEST(TestUniqueCountDispatchExplicit);
 
 template <typename ForwardIterator>
-typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(my_tag, ForwardIterator, ForwardIterator)
+typename ::cuda::std::iterator_traits<ForwardIterator>::difference_type
+unique_count(my_tag, ForwardIterator, ForwardIterator)
 {
   return 13;
 }
diff --git a/thrust/testing/unittest/assertions.h b/thrust/testing/unittest/assertions.h
index c6404d156e4..7fb924ade7f 100644
--- a/thrust/testing/unittest/assertions.h
+++ b/thrust/testing/unittest/assertions.h
@@ -409,8 +409,8 @@ void assert_equal(
   const std::string& filename = "unknown",
   int lineno                  = -1)
 {
-  using difference_type = typename THRUST_NS_QUALIFIER::iterator_difference<ForwardIterator1>::type;
-  using InputType       = typename THRUST_NS_QUALIFIER::iterator_value<ForwardIterator1>::type;
+  using difference_type = THRUST_NS_QUALIFIER::detail::it_difference_t<ForwardIterator1>;
+  using InputType       = THRUST_NS_QUALIFIER::detail::it_value_t<ForwardIterator1>;
 
   bool failure = false;
 
@@ -486,7 +486,7 @@ void assert_equal(
   const std::string& filename = "unknown",
   int lineno                  = -1)
 {
-  using InputType = typename THRUST_NS_QUALIFIER::iterator_traits<ForwardIterator1>::value_type;
+  using InputType = typename ::cuda::std::iterator_traits<ForwardIterator1>::value_type;
   assert_equal(first1, last1, first2, last2, THRUST_NS_QUALIFIER::equal_to<InputType>(), filename, lineno);
 }
 
@@ -501,7 +501,7 @@ void assert_almost_equal(
   const double a_tol          = DEFAULT_ABSOLUTE_TOL,
   const double r_tol          = DEFAULT_RELATIVE_TOL)
 {
-  using InputType = typename THRUST_NS_QUALIFIER::iterator_traits<ForwardIterator1>::value_type;
+  using InputType = typename ::cuda::std::iterator_traits<ForwardIterator1>::value_type;
   assert_equal(first1, last1, first2, last2, almost_equal_to<InputType>(a_tol, r_tol), filename, lineno);
 }
 
diff --git a/thrust/testing/zip_iterator.cu b/thrust/testing/zip_iterator.cu
index 786b839099b..aad7afc0dd8 100644
--- a/thrust/testing/zip_iterator.cu
+++ b/thrust/testing/zip_iterator.cu
@@ -109,7 +109,7 @@ struct TestZipIteratorReference
     using IteratorTuple1 = tuple<Iterator1, Iterator2>;
     using ZipIterator1   = zip_iterator<IteratorTuple1>;
 
-    using zip_iterator_reference_type1 = typename iterator_reference<ZipIterator1>::type;
+    using zip_iterator_reference_type1 = thrust::detail::it_reference_t<ZipIterator1>;
 
     host_vector<T> h_variable(1);
 
@@ -128,7 +128,7 @@ struct TestZipIteratorReference
     using IteratorTuple2 = tuple<Iterator3, Iterator4>;
     using ZipIterator2   = zip_iterator<IteratorTuple2>;
 
-    using zip_iterator_reference_type2 = typename iterator_reference<ZipIterator2>::type;
+    using zip_iterator_reference_type2 = thrust::detail::it_reference_t<ZipIterator2>;
 
     device_vector<T> d_variable(1);
 
diff --git a/thrust/thrust/advance.h b/thrust/thrust/advance.h
index 801568b6e00..404e8a31c88 100644
--- a/thrust/thrust/advance.h
+++ b/thrust/thrust/advance.h
@@ -101,7 +101,7 @@ template <typename InputIterator, typename Distance>
 _CCCL_HOST_DEVICE
 InputIterator next(
   InputIterator i
-, typename iterator_traits<InputIterator>::difference_type n = 1
+, typename ::cuda::std::iterator_traits<InputIterator>::difference_type n = 1
 );
 #endif
 
@@ -135,7 +135,7 @@ template <typename BidirectionalIterator, typename Distance>
 _CCCL_HOST_DEVICE
 BidirectionalIterator prev(
   BidirectionalIterator i
-, typename iterator_traits<BidirectionalIterator>::difference_type n = 1
+, typename ::cuda::std::iterator_traits<BidirectionalIterator>::difference_type n = 1
 );
 #endif
 
diff --git a/thrust/thrust/count.h b/thrust/thrust/count.h
index 079322802b3..aac7457e59a 100644
--- a/thrust/thrust/count.h
+++ b/thrust/thrust/count.h
@@ -89,7 +89,7 @@ THRUST_NAMESPACE_BEGIN
  *  \see https://en.cppreference.com/w/cpp/algorithm/count
  */
 template <typename DerivedPolicy, typename InputIterator, typename EqualityComparable>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
       InputIterator first,
       InputIterator last,
@@ -131,7 +131,7 @@ count(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
  *  \see https://en.cppreference.com/w/cpp/algorithm/count
  */
 template <typename InputIterator, typename EqualityComparable>
-typename thrust::iterator_traits<InputIterator>::difference_type
+thrust::detail::it_difference_t<InputIterator>
 count(InputIterator first, InputIterator last, const EqualityComparable& value);
 
 /*! \p count_if finds the number of elements in <tt>[first,last)</tt> for which
@@ -184,7 +184,7 @@ count(InputIterator first, InputIterator last, const EqualityComparable& value);
  *  \see https://en.cppreference.com/w/cpp/algorithm/count
  */
 template <typename DerivedPolicy, typename InputIterator, typename Predicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count_if(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
          InputIterator first,
          InputIterator last,
@@ -234,8 +234,7 @@ count_if(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
  *  \see https://en.cppreference.com/w/cpp/algorithm/count
  */
 template <typename InputIterator, typename Predicate>
-typename thrust::iterator_traits<InputIterator>::difference_type
-count_if(InputIterator first, InputIterator last, Predicate pred);
+thrust::detail::it_difference_t<InputIterator> count_if(InputIterator first, InputIterator last, Predicate pred);
 
 /*! \} // end counting
  *  \} // end reductions
diff --git a/thrust/thrust/detail/advance.inl b/thrust/thrust/detail/advance.inl
index fe84d373b59..ec49d604f92 100644
--- a/thrust/thrust/detail/advance.inl
+++ b/thrust/thrust/detail/advance.inl
@@ -43,7 +43,7 @@ _CCCL_HOST_DEVICE void advance(InputIterator& i, Distance n)
 }
 
 template <typename InputIterator>
-_CCCL_HOST_DEVICE InputIterator next(InputIterator i, typename iterator_traits<InputIterator>::difference_type n = 1)
+_CCCL_HOST_DEVICE InputIterator next(InputIterator i, thrust::detail::it_difference_t<InputIterator> n = 1)
 {
   thrust::system::detail::generic::advance(i, n);
   return i;
@@ -51,15 +51,16 @@ _CCCL_HOST_DEVICE InputIterator next(InputIterator i, typename iterator_traits<I
 
 template <typename BidirectionalIterator>
 _CCCL_HOST_DEVICE BidirectionalIterator
-prev(BidirectionalIterator i, typename iterator_traits<BidirectionalIterator>::difference_type n = 1)
+prev(BidirectionalIterator i, thrust::detail::it_difference_t<BidirectionalIterator> n = 1)
 {
   thrust::system::detail::generic::advance(i, -n);
   return i;
 }
 
+// FIXME(bgruber): what does this prevent against?
 template <typename BidirectionalIterator>
 _CCCL_HOST_DEVICE
-typename detail::disable_if<has_difference_type<iterator_traits<BidirectionalIterator>>::value,
+typename detail::disable_if<has_difference_type<::cuda::std::iterator_traits<BidirectionalIterator>>::value,
                             BidirectionalIterator>::type
 prev(BidirectionalIterator i, typename detail::pointer_traits<BidirectionalIterator>::difference_type n = 1)
 {
diff --git a/thrust/thrust/detail/allocator/copy_construct_range.inl b/thrust/thrust/detail/allocator/copy_construct_range.inl
index 42dd32f045e..c02f373b16e 100644
--- a/thrust/thrust/detail/allocator/copy_construct_range.inl
+++ b/thrust/thrust/detail/allocator/copy_construct_range.inl
@@ -100,12 +100,12 @@ _CCCL_HOST_DEVICE enable_if_convertible_t<FromSystem, ToSystem, Pointer> uniniti
   ZipIterator end   = begin;
 
   // get a zip_iterator pointing to the end
-  const typename thrust::iterator_difference<InputIterator>::type n = thrust::distance(first, last);
+  const thrust::detail::it_difference_t<InputIterator> n = thrust::distance(first, last);
   thrust::advance(end, n);
 
   // create a functor
-  using InputType  = typename iterator_traits<InputIterator>::value_type;
-  using OutputType = typename iterator_traits<Pointer>::value_type;
+  using InputType  = it_value_t<InputIterator>;
+  using OutputType = it_value_t<Pointer>;
 
   // do the for_each
   // note we use to_system to dispatch the for_each
@@ -135,8 +135,8 @@ _CCCL_HOST_DEVICE enable_if_convertible_t<FromSystem, ToSystem, Pointer> uniniti
   ZipIterator begin = thrust::make_zip_iterator(thrust::make_tuple(first, result));
 
   // create a functor
-  using InputType  = typename iterator_traits<InputIterator>::value_type;
-  using OutputType = typename iterator_traits<Pointer>::value_type;
+  using InputType  = it_value_t<InputIterator>;
+  using OutputType = it_value_t<Pointer>;
 
   // do the for_each_n
   // note we use to_system to dispatch the for_each_n
diff --git a/thrust/thrust/detail/allocator/tagged_allocator.h b/thrust/thrust/detail/allocator/tagged_allocator.h
index bdac004d9bb..c7cc09ea713 100644
--- a/thrust/thrust/detail/allocator/tagged_allocator.h
+++ b/thrust/thrust/detail/allocator/tagged_allocator.h
@@ -60,8 +60,8 @@ class tagged_allocator
   using value_type      = T;
   using pointer         = typename thrust::detail::pointer_traits<Pointer>::template rebind<T>::other;
   using const_pointer   = typename thrust::detail::pointer_traits<Pointer>::template rebind<const T>::other;
-  using reference       = typename thrust::iterator_reference<pointer>::type;
-  using const_reference = typename thrust::iterator_reference<const_pointer>::type;
+  using reference       = thrust::detail::it_reference_t<pointer>;
+  using const_reference = thrust::detail::it_reference_t<const_pointer>;
   using size_type       = std::size_t;
   using difference_type = typename thrust::detail::pointer_traits<pointer>::difference_type;
   using system_type     = Tag;
diff --git a/thrust/thrust/detail/count.h b/thrust/thrust/detail/count.h
index e5ce2e2b30b..ac3668190f1 100644
--- a/thrust/thrust/detail/count.h
+++ b/thrust/thrust/detail/count.h
@@ -30,26 +30,25 @@
 THRUST_NAMESPACE_BEGIN
 
 template <typename DerivedPolicy, typename InputIterator, typename EqualityComparable>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
       InputIterator first,
       InputIterator last,
       const EqualityComparable& value);
 
 template <typename DerivedPolicy, typename InputIterator, typename Predicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count_if(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
          InputIterator first,
          InputIterator last,
          Predicate pred);
 
 template <typename InputIterator, typename EqualityComparable>
-typename thrust::iterator_traits<InputIterator>::difference_type
+thrust::detail::it_difference_t<InputIterator>
 count(InputIterator first, InputIterator last, const EqualityComparable& value);
 
 template <typename InputIterator, typename Predicate>
-typename thrust::iterator_traits<InputIterator>::difference_type
-count_if(InputIterator first, InputIterator last, Predicate pred);
+thrust::detail::it_difference_t<InputIterator> count_if(InputIterator first, InputIterator last, Predicate pred);
 
 THRUST_NAMESPACE_END
 
diff --git a/thrust/thrust/detail/count.inl b/thrust/thrust/detail/count.inl
index 28876d8f12f..7813b48c28b 100644
--- a/thrust/thrust/detail/count.inl
+++ b/thrust/thrust/detail/count.inl
@@ -35,7 +35,7 @@ THRUST_NAMESPACE_BEGIN
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename DerivedPolicy, typename InputIterator, typename EqualityComparable>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
       InputIterator first,
       InputIterator last,
@@ -47,7 +47,7 @@ count(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename DerivedPolicy, typename InputIterator, typename Predicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count_if(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
          InputIterator first,
          InputIterator last,
@@ -58,7 +58,7 @@ count_if(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
 } // end count_if()
 
 template <typename InputIterator, typename EqualityComparable>
-typename thrust::iterator_traits<InputIterator>::difference_type
+thrust::detail::it_difference_t<InputIterator>
 count(InputIterator first, InputIterator last, const EqualityComparable& value)
 {
   using thrust::system::detail::generic::select_system;
@@ -71,8 +71,7 @@ count(InputIterator first, InputIterator last, const EqualityComparable& value)
 } // end count()
 
 template <typename InputIterator, typename Predicate>
-typename thrust::iterator_traits<InputIterator>::difference_type
-count_if(InputIterator first, InputIterator last, Predicate pred)
+thrust::detail::it_difference_t<InputIterator> count_if(InputIterator first, InputIterator last, Predicate pred)
 {
   using thrust::system::detail::generic::select_system;
 
diff --git a/thrust/thrust/detail/distance.inl b/thrust/thrust/detail/distance.inl
index 7316dfa91dd..b8a71b95c1d 100644
--- a/thrust/thrust/detail/distance.inl
+++ b/thrust/thrust/detail/distance.inl
@@ -34,8 +34,7 @@ THRUST_NAMESPACE_BEGIN
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename InputIterator>
-inline _CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
-distance(InputIterator first, InputIterator last)
+inline _CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator> distance(InputIterator first, InputIterator last)
 {
   return thrust::system::detail::generic::distance(first, last);
 } // end distance()
diff --git a/thrust/thrust/detail/get_iterator_value.h b/thrust/thrust/detail/get_iterator_value.h
index 330aea557cc..5779ded3c9f 100644
--- a/thrust/thrust/detail/get_iterator_value.h
+++ b/thrust/thrust/detail/get_iterator_value.h
@@ -39,8 +39,7 @@ namespace detail
 // --------------------------------------------------
 // it is okay to dereference iterator in the usual way
 template <typename DerivedPolicy, typename Iterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<Iterator>::value_type
-get_iterator_value(thrust::execution_policy<DerivedPolicy>&, Iterator it)
+_CCCL_HOST_DEVICE it_value_t<Iterator> get_iterator_value(thrust::execution_policy<DerivedPolicy>&, Iterator it)
 {
   return *it;
 } // get_iterator_value(exec,Iterator);
diff --git a/thrust/thrust/detail/memory_algorithms.h b/thrust/thrust/detail/memory_algorithms.h
index 1da74c4b56d..7ae020c52c3 100644
--- a/thrust/thrust/detail/memory_algorithms.h
+++ b/thrust/thrust/detail/memory_algorithms.h
@@ -64,7 +64,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy(ForwardIt first, ForwardIt last) noexcept
 template <typename Allocator, typename ForwardIt>
 _CCCL_HOST_DEVICE ForwardIt destroy(Allocator const& alloc, ForwardIt first, ForwardIt last) noexcept
 {
-  using T      = typename iterator_traits<ForwardIt>::value_type;
+  using T      = detail::it_value_t<ForwardIt>;
   using traits = typename detail::allocator_traits<
     ::cuda::std::remove_cv_t<::cuda::std::remove_reference_t<Allocator>>>::template rebind_traits<T>::other;
 
@@ -92,7 +92,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy_n(ForwardIt first, Size n) noexcept
 template <typename Allocator, typename ForwardIt, typename Size>
 _CCCL_HOST_DEVICE ForwardIt destroy_n(Allocator const& alloc, ForwardIt first, Size n) noexcept
 {
-  using T      = typename iterator_traits<ForwardIt>::value_type;
+  using T      = detail::it_value_t<ForwardIt>;
   using traits = typename detail::allocator_traits<
     ::cuda::std::remove_cv_t<::cuda::std::remove_reference_t<Allocator>>>::template rebind_traits<T>::other;
 
@@ -109,7 +109,7 @@ _CCCL_HOST_DEVICE ForwardIt destroy_n(Allocator const& alloc, ForwardIt first, S
 template <typename ForwardIt, typename... Args>
 _CCCL_HOST_DEVICE void uninitialized_construct(ForwardIt first, ForwardIt last, Args const&... args)
 {
-  using T = typename iterator_traits<ForwardIt>::value_type;
+  using T = detail::it_value_t<ForwardIt>;
 
   ForwardIt current = first;
 
@@ -132,7 +132,7 @@ _CCCL_HOST_DEVICE void uninitialized_construct(ForwardIt first, ForwardIt last,
 template <typename Allocator, typename ForwardIt, typename... Args>
 void uninitialized_construct_with_allocator(Allocator const& alloc, ForwardIt first, ForwardIt last, Args const&... args)
 {
-  using T      = typename iterator_traits<ForwardIt>::value_type;
+  using T      = detail::it_value_t<ForwardIt>;
   using traits = typename detail::allocator_traits<
     typename std::remove_cv<typename std::remove_reference<Allocator>::type>::type>::template rebind_traits<T>;
 
@@ -159,7 +159,7 @@ void uninitialized_construct_with_allocator(Allocator const& alloc, ForwardIt fi
 template <typename ForwardIt, typename Size, typename... Args>
 void uninitialized_construct_n(ForwardIt first, Size n, Args const&... args)
 {
-  using T = typename iterator_traits<ForwardIt>::value_type;
+  using T = detail::it_value_t<ForwardIt>;
 
   ForwardIt current = first;
 
@@ -182,7 +182,7 @@ void uninitialized_construct_n(ForwardIt first, Size n, Args const&... args)
 template <typename Allocator, typename ForwardIt, typename Size, typename... Args>
 void uninitialized_construct_n_with_allocator(Allocator const& alloc, ForwardIt first, Size n, Args const&... args)
 {
-  using T      = typename iterator_traits<ForwardIt>::value_type;
+  using T      = detail::it_value_t<ForwardIt>;
   using traits = typename detail::allocator_traits<
     typename std::remove_cv<typename std::remove_reference<Allocator>::type>::type>::template rebind_traits<T>;
 
diff --git a/thrust/thrust/detail/overlapped_copy.h b/thrust/thrust/detail/overlapped_copy.h
index f5b11525748..98e49995215 100644
--- a/thrust/thrust/detail/overlapped_copy.h
+++ b/thrust/thrust/detail/overlapped_copy.h
@@ -95,7 +95,7 @@ RandomAccessIterator2 overlapped_copy(
   RandomAccessIterator1 last,
   RandomAccessIterator2 result)
 {
-  using value_type = typename thrust::iterator_value<RandomAccessIterator1>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator1>;
 
   // make a temporary copy of [first,last), and copy into it first
   thrust::detail::temporary_array<value_type, DerivedPolicy> temp(exec, first, last);
diff --git a/thrust/thrust/detail/pointer.h b/thrust/thrust/detail/pointer.h
index b9f2fabc244..43b38f5bf76 100644
--- a/thrust/thrust/detail/pointer.h
+++ b/thrust/thrust/detail/pointer.h
@@ -1,3 +1,4 @@
+
 /*
  *  Copyright 2008-2021 NVIDIA Corporation
  *
@@ -43,29 +44,15 @@
 #include <ostream>
 
 THRUST_NAMESPACE_BEGIN
-
 template <typename Element, typename Tag, typename Reference = use_default, typename Derived = use_default>
 class pointer;
-
-// Specialize `thrust::iterator_traits` to avoid problems with the name of
-// pointer's constructor shadowing its nested pointer type. We do this before
-// pointer is defined so the specialization is correctly used inside the
-// definition.
-template <typename Element, typename Tag, typename Reference, typename Derived>
-struct iterator_traits<thrust::pointer<Element, Tag, Reference, Derived>>
-{
-  using pointer           = thrust::pointer<Element, Tag, Reference, Derived>;
-  using iterator_category = typename pointer::iterator_category;
-  using value_type        = typename pointer::value_type;
-  using difference_type   = typename pointer::difference_type;
-  using reference         = typename pointer::reference;
-};
-
 THRUST_NAMESPACE_END
 
+// Specialize `std::iterator_traits` (picked up by cuda::std::iterator_traits) to avoid problems with the name of
+// pointer's constructor shadowing its nested pointer type. We do this before pointer is defined so the specialization
+// is correctly used inside the definition.
 namespace std
 {
-
 template <typename Element, typename Tag, typename Reference, typename Derived>
 struct iterator_traits<THRUST_NS_QUALIFIER::pointer<Element, Tag, Reference, Derived>>
 {
@@ -75,7 +62,6 @@ struct iterator_traits<THRUST_NS_QUALIFIER::pointer<Element, Tag, Reference, Der
   using difference_type   = typename pointer::difference_type;
   using reference         = typename pointer::reference;
 };
-
 } // namespace std
 
 THRUST_NAMESPACE_BEGIN
diff --git a/thrust/thrust/detail/range/head_flags.h b/thrust/thrust/detail/range/head_flags.h
index aed28a5a7d3..bef88ff6aeb 100644
--- a/thrust/thrust/detail/range/head_flags.h
+++ b/thrust/thrust/detail/range/head_flags.h
@@ -36,9 +36,9 @@ namespace detail
 {
 
 template <typename RandomAccessIterator,
-          typename BinaryPredicate = thrust::equal_to<typename thrust::iterator_value<RandomAccessIterator>::type>,
+          typename BinaryPredicate = equal_to<it_value_t<RandomAccessIterator>>,
           typename ValueType       = bool,
-          typename IndexType       = typename thrust::iterator_difference<RandomAccessIterator>::type>
+          typename IndexType       = it_difference_t<RandomAccessIterator>>
 class head_flags
 {
 public:
diff --git a/thrust/thrust/detail/range/tail_flags.h b/thrust/thrust/detail/range/tail_flags.h
index 48470c0f369..456bd2f33e3 100644
--- a/thrust/thrust/detail/range/tail_flags.h
+++ b/thrust/thrust/detail/range/tail_flags.h
@@ -36,9 +36,9 @@ namespace detail
 {
 
 template <typename RandomAccessIterator,
-          typename BinaryPredicate = thrust::equal_to<typename thrust::iterator_value<RandomAccessIterator>::type>,
+          typename BinaryPredicate = equal_to<detail::it_value_t<RandomAccessIterator>>,
           typename ValueType       = bool,
-          typename IndexType       = typename thrust::iterator_difference<RandomAccessIterator>::type>
+          typename IndexType       = detail::it_difference_t<RandomAccessIterator>>
 class tail_flags
 {
   // XXX WAR cudafe bug
diff --git a/thrust/thrust/detail/reduce.inl b/thrust/thrust/detail/reduce.inl
index 2d7f7cd4729..c008ec17e6e 100644
--- a/thrust/thrust/detail/reduce.inl
+++ b/thrust/thrust/detail/reduce.inl
@@ -38,7 +38,7 @@ THRUST_NAMESPACE_BEGIN
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename DerivedPolicy, typename InputIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::value_type
+_CCCL_HOST_DEVICE detail::it_value_t<InputIterator>
 reduce(const thrust::detail::execution_policy_base<DerivedPolicy>& exec, InputIterator first, InputIterator last)
 {
   using thrust::system::detail::generic::reduce;
@@ -149,7 +149,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> reduce_by_key(
 } // end reduce_by_key()
 
 template <typename InputIterator>
-typename thrust::iterator_traits<InputIterator>::value_type reduce(InputIterator first, InputIterator last)
+detail::it_value_t<InputIterator> reduce(InputIterator first, InputIterator last)
 {
   using thrust::system::detail::generic::select_system;
 
diff --git a/thrust/thrust/detail/temporary_array.h b/thrust/thrust/detail/temporary_array.h
index 848255aac89..4fee9053101 100644
--- a/thrust/thrust/detail/temporary_array.h
+++ b/thrust/thrust/detail/temporary_array.h
@@ -128,7 +128,7 @@ template <typename Iterator, typename FromSystem, typename ToSystem>
 struct move_to_system_base
     : public eval_if<::cuda::std::is_convertible<FromSystem, ToSystem>::value,
                      identity_<tagged_iterator_range<Iterator, ToSystem>>,
-                     identity_<temporary_array<typename thrust::iterator_value<Iterator>::type, ToSystem>>>
+                     identity_<temporary_array<thrust::detail::it_value_t<Iterator>, ToSystem>>>
 {};
 
 template <typename Iterator, typename FromSystem, typename ToSystem>
diff --git a/thrust/thrust/detail/trivial_sequence.h b/thrust/thrust/detail/trivial_sequence.h
index 85f51eadfd4..4ae2ef481b7 100644
--- a/thrust/thrust/detail/trivial_sequence.h
+++ b/thrust/thrust/detail/trivial_sequence.h
@@ -85,7 +85,7 @@ struct _trivial_sequence<Iterator, DerivedPolicy, thrust::detail::true_type>
 template <typename Iterator, typename DerivedPolicy>
 struct _trivial_sequence<Iterator, DerivedPolicy, thrust::detail::false_type>
 {
-  using iterator_value = typename thrust::iterator_value<Iterator>::type;
+  using iterator_value = it_value_t<Iterator>;
   using iterator_type  = typename thrust::detail::temporary_array<iterator_value, DerivedPolicy>::iterator;
 
   thrust::detail::temporary_array<iterator_value, DerivedPolicy> buffer;
diff --git a/thrust/thrust/detail/type_traits/iterator/is_output_iterator.h b/thrust/thrust/detail/type_traits/iterator/is_output_iterator.h
index 2741ca38212..1d06146955f 100644
--- a/thrust/thrust/detail/type_traits/iterator/is_output_iterator.h
+++ b/thrust/thrust/detail/type_traits/iterator/is_output_iterator.h
@@ -25,38 +25,22 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
-#include <thrust/detail/type_traits.h>
-#include <thrust/detail/type_traits/is_metafunction_defined.h>
+
 #include <thrust/iterator/detail/any_assign.h>
 #include <thrust/iterator/iterator_traits.h>
 
+#include <cuda/std/type_traits>
+
 THRUST_NAMESPACE_BEGIN
 
 namespace detail
 {
+template <typename T, typename SFINAE = void>
+_CCCL_INLINE_VAR constexpr bool is_output_iterator = true;
 
 template <typename T>
-struct is_void_like
-    : ::cuda::std::disjunction<::cuda::std::is_void<T>, ::cuda::std::is_same<T, thrust::detail::any_assign>>
-{}; // end is_void_like
-
-template <typename T>
-struct lazy_is_void_like : is_void_like<typename T::type>
-{}; // end lazy_is_void_like
-
-// XXX this meta function should first check that T is actually an iterator
-//
-//     if thrust::iterator_value<T> is defined and thrust::iterator_value<T>::type == void
-//       return false
-//     else
-//       return true
-template <typename T>
-struct is_output_iterator
-    : eval_if<is_metafunction_defined<thrust::iterator_value<T>>::value,
-              lazy_is_void_like<thrust::iterator_value<T>>,
-              thrust::detail::true_type>::type
-{}; // end is_output_iterator
-
+_CCCL_INLINE_VAR constexpr bool is_output_iterator<T, ::cuda::std::void_t<it_value_t<T>>> =
+  ::cuda::std::is_void_v<it_value_t<T>> || ::cuda::std::is_same_v<it_value_t<T>, any_assign>;
 } // namespace detail
 
 THRUST_NAMESPACE_END
diff --git a/thrust/thrust/detail/unique.inl b/thrust/thrust/detail/unique.inl
index 2750cb5e10d..2eec8dba58f 100644
--- a/thrust/thrust/detail/unique.inl
+++ b/thrust/thrust/detail/unique.inl
@@ -303,7 +303,7 @@ thrust::pair<OutputIterator1, OutputIterator2> unique_by_key_copy(
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator> unique_count(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
   ForwardIterator first,
   ForwardIterator last,
@@ -315,7 +315,7 @@ _CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename DerivedPolicy, typename ForwardIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator> unique_count(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last)
 {
   using thrust::system::detail::generic::unique_count;
@@ -324,7 +324,7 @@ _CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename ForwardIterator, typename BinaryPredicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator>
 unique_count(ForwardIterator first, ForwardIterator last, BinaryPredicate binary_pred)
 {
   using thrust::system::detail::generic::select_system;
@@ -338,7 +338,7 @@ unique_count(ForwardIterator first, ForwardIterator last, BinaryPredicate binary
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename ForwardIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator>
 unique_count(ForwardIterator first, ForwardIterator last)
 {
   using thrust::system::detail::generic::select_system;
diff --git a/thrust/thrust/detail/vector_base.inl b/thrust/thrust/detail/vector_base.inl
index d66765775ce..9109dd8f502 100644
--- a/thrust/thrust/detail/vector_base.inl
+++ b/thrust/thrust/detail/vector_base.inl
@@ -1117,7 +1117,7 @@ bool vector_equal(InputIterator1 first1, InputIterator1 last1, InputIterator2 fi
 template <typename InputIterator1, typename InputIterator2>
 bool vector_equal(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, thrust::detail::false_type)
 {
-  typename thrust::iterator_difference<InputIterator1>::type n = thrust::distance(first1, last1);
+  it_difference_t<InputIterator1> n = thrust::distance(first1, last1);
 
   using FromSystem1 = typename thrust::iterator_system<InputIterator1>::type;
   using FromSystem2 = typename thrust::iterator_system<InputIterator2>::type;
diff --git a/thrust/thrust/distance.h b/thrust/thrust/distance.h
index 81ca59da012..1cbff47cffe 100644
--- a/thrust/thrust/distance.h
+++ b/thrust/thrust/distance.h
@@ -70,7 +70,7 @@ THRUST_NAMESPACE_BEGIN
  *  \see https://en.cppreference.com/w/cpp/iterator/distance
  */
 template <typename InputIterator>
-inline _CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+inline _CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 distance(InputIterator first, InputIterator last);
 
 /*! \} // end iterators
diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h
index b29e9cbdb5f..4f159fa5432 100644
--- a/thrust/thrust/iterator/counting_iterator.h
+++ b/thrust/thrust/iterator/counting_iterator.h
@@ -128,7 +128,7 @@ struct make_counting_iterator_base
                                    eval_if<::cuda::std::is_integral<Incrementable>::value,
                                            numeric_difference<Incrementable>,
                                            identity_<::cuda::std::ptrdiff_t>>,
-                                   iterator_difference<Incrementable>>>;
+                                   lazy_trait<it_difference_t, Incrementable>>>;
 
   // our implementation departs from Boost's in that counting_iterator::dereference
   // returns a copy of its counter, rather than a reference to it. returning a reference
diff --git a/thrust/thrust/iterator/detail/iterator_adaptor_base.h b/thrust/thrust/iterator/detail/iterator_adaptor_base.h
index 5e42c9edd59..32452490741 100644
--- a/thrust/thrust/iterator/detail/iterator_adaptor_base.h
+++ b/thrust/thrust/iterator/detail/iterator_adaptor_base.h
@@ -61,15 +61,15 @@ template <typename Derived,
 struct make_iterator_adaptor_base
 {
 private:
-  using value     = replace_if_use_default<Value, iterator_value<Base>>;
+  using value     = replace_if_use_default<Value, lazy_trait<it_value_t, Base>>;
   using system    = replace_if_use_default<System, iterator_system<Base>>;
   using traversal = replace_if_use_default<Traversal, iterator_traversal<Base>>;
   using reference =
     replace_if_use_default<Reference,
                            ::cuda::std::_If<::cuda::std::is_same_v<Value, use_default>,
-                                            iterator_reference<Base>,
+                                            lazy_trait<it_reference_t, Base>,
                                             ::cuda::std::add_lvalue_reference<Value>>>;
-  using difference = replace_if_use_default<Difference, iterator_difference<Base>>;
+  using difference = replace_if_use_default<Difference, lazy_trait<it_difference_t, Base>>;
 
 public:
   using type = iterator_facade<Derived, value, system, traversal, reference, difference>;
diff --git a/thrust/thrust/iterator/detail/tagged_iterator.h b/thrust/thrust/iterator/detail/tagged_iterator.h
index 1c21fbaca65..a97f99e5de1 100644
--- a/thrust/thrust/iterator/detail/tagged_iterator.h
+++ b/thrust/thrust/iterator/detail/tagged_iterator.h
@@ -40,11 +40,11 @@ template <typename Iterator, typename Tag>
 using make_tagged_iterator_base =
   iterator_adaptor<tagged_iterator<Iterator, Tag>,
                    Iterator,
-                   typename iterator_value<Iterator>::type,
+                   it_value_t<Iterator>,
                    Tag,
                    typename iterator_traversal<Iterator>::type,
-                   typename iterator_reference<Iterator>::type,
-                   typename iterator_difference<Iterator>::type>;
+                   it_reference_t<Iterator>,
+                   it_difference_t<Iterator>>;
 
 template <typename Iterator, typename Tag>
 class tagged_iterator : public make_tagged_iterator_base<Iterator, Tag>
diff --git a/thrust/thrust/iterator/iterator_traits.h b/thrust/thrust/iterator/iterator_traits.h
index 225f1a173c4..d8c0f6b17a2 100644
--- a/thrust/thrust/iterator/iterator_traits.h
+++ b/thrust/thrust/iterator/iterator_traits.h
@@ -45,68 +45,103 @@
 #include <thrust/iterator/iterator_categories.h>
 
 #include <cuda/std/__type_traits/void_t.h>
-
-#if _CCCL_COMPILER(NVRTC)
-#  include <cuda/std/iterator>
-#else // _CCCL_COMPILER(NVRTC)
-#  include <iterator>
-#endif // _CCCL_COMPILER(NVRTC)
+#include <cuda/std/iterator>
 
 THRUST_NAMESPACE_BEGIN
 
+namespace detail
+{
+// the following iterator helpers are not named it_value_t etc, like the C++20 facilities, because they are defined in
+// terms of C++17 iterator_traits and not the new C++20 indirectly_readable trait etc. This allows them to detect nested
+// value_type, difference_type and reference aliases, which the new C+20 traits do not consider (they only consider
+// specializations of iterator_traits). Also, a value_type of void remains supported (needed by some output iterators).
+
+template <typename It>
+using it_value_t = typename ::cuda::std::iterator_traits<It>::value_type;
+
+template <typename It>
+using it_reference_t = typename ::cuda::std::iterator_traits<It>::reference;
+
+template <typename It>
+using it_difference_t = typename ::cuda::std::iterator_traits<It>::difference_type;
+
+template <typename It>
+using it_pointer_t = typename ::cuda::std::iterator_traits<It>::pointer;
+
+// use this whenever you need to lazily evaluate a trait. E.g., as an alternative in replace_if_use_default.
+template <template <typename...> typename Trait, typename... Args>
+struct lazy_trait
+{
+  using type = Trait<Args...>;
+};
+} // namespace detail
+
 //! \p iterator_traits is a type trait class that provides a uniform interface for querying the properties of iterators
-//! at compile-time.
+//! at compile-time. You can specialize cuda::std::iterator_traits for your own iterator types if needed.
+//! deprecated [Since 3.0]
 template <typename T>
-struct iterator_traits
-    :
+using iterator_traits CCCL_DEPRECATED_BECAUSE("Use cuda::std::iterator_traits instead") =
+// FIXME(bgruber): switching to ::cuda::std::iterator_traits<T> breaks some tests, e.g. cub.test.device_merge_sort.lid_1
 #if _CCCL_COMPILER(NVRTC)
-    ::cuda
+  ::cuda
 #endif // _CCCL_COMPILER(NVRTC)
-    ::std::iterator_traits<T>
-{};
+  ::std::iterator_traits<T>;
+
+_CCCL_SUPPRESS_DEPRECATED_PUSH
 
 // value
 
+//! deprecated [Since 3.0]
 template <typename Iterator>
-struct iterator_value
+struct CCCL_DEPRECATED_BECAUSE("Use cuda::std::it_value_t instead") iterator_value
 {
   using type = typename iterator_traits<Iterator>::value_type;
 };
 
+//! deprecated [Since 3.0]
 template <typename Iterator>
-using iterator_value_t = typename iterator_value<Iterator>::type;
+using iterator_value_t CCCL_DEPRECATED_BECAUSE("Use cuda::std::it_value_t instead") = iterator_value<Iterator>;
 
 // pointer
 
+//! deprecated [Since 3.0]
 template <typename Iterator>
-struct iterator_pointer
+struct CCCL_DEPRECATED iterator_pointer
 {
   using type = typename iterator_traits<Iterator>::pointer;
 };
+
+//! deprecated [Since 3.0]
 template <typename Iterator>
-using iterator_pointer_t = typename iterator_pointer<Iterator>::type;
+using iterator_pointer_t CCCL_DEPRECATED = typename iterator_pointer<Iterator>::type;
 
 // reference
 
+//! deprecated [Since 3.0]
 template <typename Iterator>
-struct iterator_reference
+struct CCCL_DEPRECATED_BECAUSE("Use cuda::std::it_reference_t instead") iterator_reference
 {
   using type = typename iterator_traits<Iterator>::reference;
 };
 
+//! deprecated [Since 3.0]
 template <typename Iterator>
-using iterator_reference_t = typename iterator_reference<Iterator>::type;
+using iterator_reference_t
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::it_reference_t instead") = typename iterator_reference<Iterator>::type;
 
 // difference
 
+//! deprecated [Since 3.0]
 template <typename Iterator>
-struct iterator_difference
+struct CCCL_DEPRECATED_BECAUSE("Use cuda::std::it_difference_t instead") iterator_difference
 {
   using type = typename iterator_traits<Iterator>::difference_type;
 };
 
+//! deprecated [Since 3.0]
 template <typename Iterator>
-using iterator_difference_t = typename iterator_difference<Iterator>::type;
+using iterator_difference_t
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::it_difference_t instead") = typename iterator_difference<Iterator>::type;
 
 // traversal
 
@@ -132,6 +167,8 @@ struct iterator_system_impl<Iterator, ::cuda::std::void_t<typename iterator_trai
 {};
 } // namespace detail
 
+_CCCL_SUPPRESS_DEPRECATED_POP
+
 template <typename Iterator>
 struct iterator_system : detail::iterator_system_impl<Iterator>
 {};
diff --git a/thrust/thrust/iterator/permutation_iterator.h b/thrust/thrust/iterator/permutation_iterator.h
index 2c7e2a69c19..85c1fb1761b 100644
--- a/thrust/thrust/iterator/permutation_iterator.h
+++ b/thrust/thrust/iterator/permutation_iterator.h
@@ -62,10 +62,10 @@ struct make_permutation_iterator_base
   using type =
     iterator_adaptor<permutation_iterator<ElementIterator, IndexIterator>,
                      IndexIterator,
-                     typename iterator_value<ElementIterator>::type,
+                     it_value_t<ElementIterator>,
                      typename minimum_system<System1, System2>::type,
                      use_default,
-                     typename iterator_reference<ElementIterator>::type>;
+                     it_reference_t<ElementIterator>>;
 };
 } // namespace detail
 
diff --git a/thrust/thrust/iterator/transform_input_output_iterator.h b/thrust/thrust/iterator/transform_input_output_iterator.h
index 0bd63b6a774..a6a5f99f8cc 100644
--- a/thrust/thrust/iterator/transform_input_output_iterator.h
+++ b/thrust/thrust/iterator/transform_input_output_iterator.h
@@ -45,7 +45,7 @@ namespace detail
 template <typename InputFunction, typename OutputFunction, typename Iterator>
 class transform_input_output_iterator_proxy
 {
-  using iterator_value_type = typename iterator_value<Iterator>::type;
+  using iterator_value_type = it_value_t<Iterator>;
   using Value               = invoke_result_t<InputFunction, iterator_value_type>;
 
 public:
@@ -90,7 +90,7 @@ template <typename InputFunction, typename OutputFunction, typename Iterator>
 struct make_transform_input_output_iterator_base
 {
 private:
-  using iterator_value_type = typename iterator_value<Iterator>::type;
+  using iterator_value_type = it_value_t<Iterator>;
 
 public:
   using type =
diff --git a/thrust/thrust/iterator/transform_iterator.h b/thrust/thrust/iterator/transform_iterator.h
index 17008d21d55..7783d39bfa6 100644
--- a/thrust/thrust/iterator/transform_iterator.h
+++ b/thrust/thrust/iterator/transform_iterator.h
@@ -63,7 +63,7 @@ template <class UnaryFunc, class Iterator>
 struct transform_iterator_reference
 {
   // by default, dereferencing the iterator yields the same as the function.
-  using type = decltype(::cuda::std::declval<UnaryFunc>()(::cuda::std::declval<iterator_value_t<Iterator>>()));
+  using type = decltype(::cuda::std::declval<UnaryFunc>()(::cuda::std::declval<it_value_t<Iterator>>()));
 };
 
 // for certain function objects, we need to tweak the reference type. Notably, identity functions must decay to values.
@@ -71,13 +71,13 @@ struct transform_iterator_reference
 template <class Iterator>
 struct transform_iterator_reference<::cuda::std::identity, Iterator>
 {
-  using type = iterator_value_t<Iterator>;
+  using type = it_value_t<Iterator>;
 };
 template <typename Eval, class Iterator>
 struct transform_iterator_reference<functional::actor<Eval>, Iterator>
 {
   using type = ::cuda::std::remove_reference_t<decltype(::cuda::std::declval<functional::actor<Eval>>()(
-    ::cuda::std::declval<iterator_value_t<Iterator>>()))>;
+    ::cuda::std::declval<it_value_t<Iterator>>()))>;
 };
 
 // Type function to compute the iterator_adaptor instantiation to be used for transform_iterator
@@ -94,7 +94,7 @@ struct make_transform_iterator_base
                      Iterator,
                      value_type,
                      use_default,
-                     typename iterator_traits<Iterator>::iterator_category,
+                     typename ::cuda::std::iterator_traits<Iterator>::iterator_category,
                      reference>;
 };
 
@@ -329,7 +329,7 @@ class transform_iterator
     // The workaround is to create a temporary to allow iterators with wrapped/proxy references to convert to their
     // value type before calling m_f. This also loads values from a different memory space (cf. `device_reference`).
     // Note that this disallows mutable operations through m_f.
-    iterator_value_t<Iterator> const& x = *this->base();
+    detail::it_value_t<Iterator> const& x = *this->base();
     // FIXME(bgruber): x may be a reference to a temporary (e.g. if the base iterator is a counting_iterator). If `m_f`
     // does not produce an independent copy and super_t::reference is a reference, we return a dangling reference (e.g.
     // for any `[thrust|::cuda::std]::identity` functor).
diff --git a/thrust/thrust/iterator/zip_iterator.h b/thrust/thrust/iterator/zip_iterator.h
index 6eecd98a152..d24d20cd15d 100644
--- a/thrust/thrust/iterator/zip_iterator.h
+++ b/thrust/thrust/iterator/zip_iterator.h
@@ -66,14 +66,14 @@ struct make_zip_iterator_base
 template <typename... Its>
 struct make_zip_iterator_base<::cuda::std::tuple<Its...>>
 {
-  // reference type is the type of the tuple obtained from the iterators' reference types.
-  using reference = tuple_of_iterator_references<iterator_reference_t<Its>...>;
+  // reference type is the type of the tuple obtained from the iterator's reference types.
+  using reference = tuple_of_iterator_references<it_reference_t<Its>...>;
 
   // Boost's Value type is the same as reference type. using value_type = reference;
-  using value_type = ::cuda::std::tuple<iterator_value_t<Its>...>;
+  using value_type = ::cuda::std::tuple<it_value_t<Its>...>;
 
   // Difference type is the first iterator's difference type
-  using difference_type = iterator_difference_t<::cuda::std::tuple_element_t<0, ::cuda::std::tuple<Its...>>>;
+  using difference_type = it_difference_t<::cuda::std::tuple_element_t<0, ::cuda::std::tuple<Its...>>>;
 
   // Iterator system is the minimum system tag in the iterator tuple
   using system = ::cuda::std::__type_fold_left<::cuda::std::__type_list<iterator_system_t<Its>...>,
@@ -324,3 +324,20 @@ inline _CCCL_HOST_DEVICE zip_iterator<_CUDA_VSTD::tuple<Iterators...>> make_zip_
 //! \} // end iterators
 
 THRUST_NAMESPACE_END
+
+// libcu++ iterator traits fail for complex zip_iterators in C++17, see e.g.: https://godbolt.org/z/7jb4qG3bb
+// The reason is that libcu++ backported the C++20 range iterator machinery to C++17, but C++17 has slightly different
+// language rules, especially regarding `void`. We deemed to it too hard to work around the issues.
+#if _CCCL_STD_VER < 2020
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+template <typename IteratorTuple>
+struct iterator_traits<THRUST_NS_QUALIFIER::zip_iterator<IteratorTuple>>
+{
+  using It                = THRUST_NS_QUALIFIER::zip_iterator<IteratorTuple>;
+  using value_type        = typename It::value_type;
+  using reference         = typename It::reference;
+  using iterator_category = typename It::iterator_category;
+  using difference_type   = typename It::difference_type;
+};
+_LIBCUDACXX_END_NAMESPACE_STD
+#endif // _CCCL_STD_VER < 2020
diff --git a/thrust/thrust/reduce.h b/thrust/thrust/reduce.h
index 49ff978bbf5..c9e0201549f 100644
--- a/thrust/thrust/reduce.h
+++ b/thrust/thrust/reduce.h
@@ -83,7 +83,7 @@ THRUST_NAMESPACE_BEGIN
  *  \see https://en.cppreference.com/w/cpp/algorithm/accumulate
  */
 template <typename DerivedPolicy, typename InputIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::value_type
+_CCCL_HOST_DEVICE detail::it_value_t<InputIterator>
 reduce(const thrust::detail::execution_policy_base<DerivedPolicy>& exec, InputIterator first, InputIterator last);
 
 /*! \p reduce is a generalization of summation: it computes the sum (or some
@@ -125,7 +125,7 @@ reduce(const thrust::detail::execution_policy_base<DerivedPolicy>& exec, InputIt
  *  \see https://en.cppreference.com/w/cpp/algorithm/accumulate
  */
 template <typename InputIterator>
-typename thrust::iterator_traits<InputIterator>::value_type reduce(InputIterator first, InputIterator last);
+detail::it_value_t<InputIterator> reduce(InputIterator first, InputIterator last);
 
 /*! \p reduce is a generalization of summation: it computes the sum (or some
  *  other binary operation) of all the elements in the range <tt>[first,
diff --git a/thrust/thrust/system/cuda/detail/adjacent_difference.h b/thrust/thrust/system/cuda/detail/adjacent_difference.h
index c62aadfc14d..13061c5b86a 100644
--- a/thrust/thrust/system/cuda/detail/adjacent_difference.h
+++ b/thrust/thrust/system/cuda/detail/adjacent_difference.h
@@ -150,8 +150,8 @@ adjacent_difference(execution_policy<Derived>& policy, InputIt first, InputIt la
   using UnwrapInputIt  = thrust::try_unwrap_contiguous_iterator_t<InputIt>;
   using UnwrapOutputIt = thrust::try_unwrap_contiguous_iterator_t<OutputIt>;
 
-  using InputValueT  = thrust::iterator_value_t<UnwrapInputIt>;
-  using OutputValueT = thrust::iterator_value_t<UnwrapOutputIt>;
+  using InputValueT  = thrust::detail::it_value_t<UnwrapInputIt>;
+  using OutputValueT = thrust::detail::it_value_t<UnwrapOutputIt>;
 
   constexpr bool can_compare_iterators =
     ::cuda::std::is_pointer<UnwrapInputIt>::value && ::cuda::std::is_pointer<UnwrapOutputIt>::value
@@ -207,7 +207,7 @@ template <class Derived, class InputIt, class OutputIt>
 OutputIt _CCCL_HOST_DEVICE
 adjacent_difference(execution_policy<Derived>& policy, InputIt first, InputIt last, OutputIt result)
 {
-  using input_type = typename iterator_traits<InputIt>::value_type;
+  using input_type = thrust::detail::it_value_t<InputIt>;
   return cuda_cub::adjacent_difference(policy, first, last, result, minus<input_type>());
 }
 
diff --git a/thrust/thrust/system/cuda/detail/async/inclusive_scan.h b/thrust/thrust/system/cuda/detail/async/inclusive_scan.h
index 2748185c3ad..826729751c7 100644
--- a/thrust/thrust/system/cuda/detail/async/inclusive_scan.h
+++ b/thrust/thrust/system/cuda/detail/async/inclusive_scan.h
@@ -66,7 +66,7 @@ template <typename DerivedPolicy, typename ForwardIt, typename Size, typename Ou
 unique_eager_event
 async_inclusive_scan_n(execution_policy<DerivedPolicy>& policy, ForwardIt first, Size n, OutputIt out, BinaryOp op)
 {
-  using AccumT     = typename thrust::iterator_traits<ForwardIt>::value_type;
+  using AccumT     = thrust::detail::it_value_t<ForwardIt>;
   using Dispatch32 = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, cub::NullType, std::uint32_t, AccumT>;
   using Dispatch64 = cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, cub::NullType, std::uint64_t, AccumT>;
 
@@ -132,8 +132,8 @@ unique_eager_event async_inclusive_scan_n(
   execution_policy<DerivedPolicy>& policy, ForwardIt first, Size n, OutputIt out, InitialValueType init, BinaryOp op)
 {
   using InputValueT = cub::detail::InputValue<InitialValueType>;
-  using AccumT      = typename ::cuda::std::
-    __accumulator_t<BinaryOp, typename ::cuda::std::iterator_traits<ForwardIt>::value_type, InitialValueType>;
+  using AccumT =
+    typename ::cuda::std::__accumulator_t<BinaryOp, thrust::detail::it_value_t<ForwardIt>, InitialValueType>;
 
   using Dispatch32 =
     cub::DispatchScan<ForwardIt, OutputIt, BinaryOp, InputValueT, std::int32_t, AccumT, cub::ForceInclusive::Yes>;
diff --git a/thrust/thrust/system/cuda/detail/copy_if.h b/thrust/thrust/system/cuda/detail/copy_if.h
index 33ba756eafa..56560e211ff 100644
--- a/thrust/thrust/system/cuda/detail/copy_if.h
+++ b/thrust/thrust/system/cuda/detail/copy_if.h
@@ -188,7 +188,7 @@ THRUST_RUNTIME_FUNCTION OutputIt copy_if(
   OutputIt output,
   Predicate predicate)
 {
-  using size_type = typename iterator_traits<InputIt>::difference_type;
+  using size_type = thrust::detail::it_difference_t<InputIt>;
 
   size_type num_items       = static_cast<size_type>(thrust::distance(first, last));
   cudaError_t status        = cudaSuccess;
diff --git a/thrust/thrust/system/cuda/detail/core/load_iterator.h b/thrust/thrust/system/cuda/detail/core/load_iterator.h
index 6f2c118b151..cf3a114c1f9 100644
--- a/thrust/thrust/system/cuda/detail/core/load_iterator.h
+++ b/thrust/thrust/system/cuda/detail/core/load_iterator.h
@@ -44,8 +44,8 @@ namespace cuda_cub::core::detail
 template <class PtxPlan, class It>
 struct LoadIterator
 {
-  using value_type = typename ::cuda::std::iterator_traits<It>::value_type;
-  using size_type  = typename ::cuda::std::iterator_traits<It>::difference_type;
+  using value_type = thrust::detail::it_value_t<It>;
+  using size_type  = thrust::detail::it_difference_t<It>;
 
   using type =
     ::cuda::std::conditional_t<is_contiguous_iterator_v<It>,
diff --git a/thrust/thrust/system/cuda/detail/core/util.h b/thrust/thrust/system/cuda/detail/core/util.h
index a5f87026d7d..c860fc56bb1 100644
--- a/thrust/thrust/system/cuda/detail/core/util.h
+++ b/thrust/thrust/system/cuda/detail/core/util.h
@@ -485,7 +485,7 @@ struct get_arch<Plan<Arch>>
 // BlockLoad
 // -----------
 // a helper metaprogram that returns type of a block loader
-template <class PtxPlan, class It, class T = typename iterator_traits<It>::value_type>
+template <class PtxPlan, class It, class T = thrust::detail::it_value_t<It>>
 struct BlockLoad
 {
   using type = cub::BlockLoad<T, PtxPlan::BLOCK_THREADS, PtxPlan::ITEMS_PER_THREAD, PtxPlan::LOAD_ALGORITHM, 1, 1>;
diff --git a/thrust/thrust/system/cuda/detail/count.h b/thrust/thrust/system/cuda/detail/count.h
index 530682ba717..22b373e18c0 100644
--- a/thrust/thrust/system/cuda/detail/count.h
+++ b/thrust/thrust/system/cuda/detail/count.h
@@ -49,10 +49,10 @@ namespace cuda_cub
 {
 
 template <class Derived, class InputIt, class UnaryPred>
-typename iterator_traits<InputIt>::difference_type _CCCL_HOST_DEVICE
+thrust::detail::it_difference_t<InputIt> _CCCL_HOST_DEVICE
 count_if(execution_policy<Derived>& policy, InputIt first, InputIt last, UnaryPred unary_pred)
 {
-  using size_type       = typename iterator_traits<InputIt>::difference_type;
+  using size_type       = thrust::detail::it_difference_t<InputIt>;
   using flag_iterator_t = transform_iterator<UnaryPred, InputIt, size_type, size_type>;
 
   return cuda_cub::reduce_n(
@@ -60,7 +60,7 @@ count_if(execution_policy<Derived>& policy, InputIt first, InputIt last, UnaryPr
 }
 
 template <class Derived, class InputIt, class Value>
-typename iterator_traits<InputIt>::difference_type _CCCL_HOST_DEVICE
+thrust::detail::it_difference_t<InputIt> _CCCL_HOST_DEVICE
 count(execution_policy<Derived>& policy, InputIt first, InputIt last, Value const& value)
 {
   return cuda_cub::count_if(policy, first, last, thrust::detail::equal_to_value<Value>(value));
diff --git a/thrust/thrust/system/cuda/detail/equal.h b/thrust/thrust/system/cuda/detail/equal.h
index dcff749eaed..a60f21dd158 100644
--- a/thrust/thrust/system/cuda/detail/equal.h
+++ b/thrust/thrust/system/cuda/detail/equal.h
@@ -55,7 +55,7 @@ equal(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1, InputI
 template <class Derived, class InputIt1, class InputIt2>
 bool _CCCL_HOST_DEVICE equal(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1, InputIt2 first2)
 {
-  using InputType1 = typename thrust::iterator_value<InputIt1>::type;
+  using InputType1 = thrust::detail::it_value_t<InputIt1>;
   return cuda_cub::equal(policy, first1, last1, first2, equal_to<InputType1>());
 }
 
diff --git a/thrust/thrust/system/cuda/detail/extrema.h b/thrust/thrust/system/cuda/detail/extrema.h
index 74891347cd1..eb9a645f00f 100644
--- a/thrust/thrust/system/cuda/detail/extrema.h
+++ b/thrust/thrust/system/cuda/detail/extrema.h
@@ -51,6 +51,8 @@
 #  include <thrust/system/cuda/detail/cdp_dispatch.h>
 #  include <thrust/system/cuda/detail/reduce.h>
 
+#  include <cuda/std/iterator>
+
 #  include <cstdint>
 
 THRUST_NAMESPACE_BEGIN
@@ -367,8 +369,8 @@ element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, BinaryPr
     return last;
   }
 
-  using InputType = typename iterator_traits<ItemsIt>::value_type;
-  using IndexType = typename iterator_traits<ItemsIt>::difference_type;
+  using InputType = thrust::detail::it_value_t<ItemsIt>;
+  using IndexType = thrust::detail::it_difference_t<ItemsIt>;
 
   IndexType num_items = static_cast<IndexType>(thrust::distance(first, last));
 
@@ -403,7 +405,7 @@ min_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, Bina
 template <class Derived, class ItemsIt>
 ItemsIt _CCCL_HOST_DEVICE min_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last)
 {
-  using value_type = typename iterator_value<ItemsIt>::type;
+  using value_type = thrust::detail::it_value_t<ItemsIt>;
   return cuda_cub::min_element(policy, first, last, less<value_type>());
 }
 
@@ -422,7 +424,7 @@ max_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, Bina
 template <class Derived, class ItemsIt>
 ItemsIt _CCCL_HOST_DEVICE max_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last)
 {
-  using value_type = typename iterator_value<ItemsIt>::type;
+  using value_type = thrust::detail::it_value_t<ItemsIt>;
   return cuda_cub::max_element(policy, first, last, less<value_type>());
 }
 
@@ -440,8 +442,7 @@ minmax_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, B
   }
 
   THRUST_CDP_DISPATCH(
-    (using InputType = typename iterator_traits<ItemsIt>::value_type;
-     using IndexType = typename iterator_traits<ItemsIt>::difference_type;
+    (using InputType = thrust::detail::it_value_t<ItemsIt>; using IndexType = thrust::detail::it_difference_t<ItemsIt>;
 
      const auto num_items = static_cast<IndexType>(thrust::distance(first, last));
 
@@ -467,7 +468,7 @@ minmax_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, B
 template <class Derived, class ItemsIt>
 pair<ItemsIt, ItemsIt> _CCCL_HOST_DEVICE minmax_element(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last)
 {
-  using value_type = typename iterator_value<ItemsIt>::type;
+  using value_type = thrust::detail::it_value_t<ItemsIt>;
   return cuda_cub::minmax_element(policy, first, last, less<value_type>());
 }
 
diff --git a/thrust/thrust/system/cuda/detail/find.h b/thrust/thrust/system/cuda/detail/find.h
index 15fb1fad993..da2a2ffdce2 100644
--- a/thrust/thrust/system/cuda/detail/find.h
+++ b/thrust/thrust/system/cuda/detail/find.h
@@ -97,7 +97,7 @@ template <class ValueType, class InputIt, class UnaryOp>
 struct transform_input_iterator_t
 {
   using self_t            = transform_input_iterator_t;
-  using difference_type   = typename iterator_traits<InputIt>::difference_type;
+  using difference_type   = thrust::detail::it_difference_t<InputIt>;
   using value_type        = ValueType;
   using pointer           = void;
   using reference         = value_type;
@@ -136,13 +136,13 @@ struct transform_input_iterator_t
 
   _CCCL_HOST_DEVICE _CCCL_FORCEINLINE reference operator*() const
   {
-    typename thrust::iterator_value<InputIt>::type x = *input;
+    thrust::detail::it_value_t<InputIt> x = *input;
     return op(x);
   }
 
   _CCCL_HOST_DEVICE _CCCL_FORCEINLINE reference operator*()
   {
-    typename thrust::iterator_value<InputIt>::type x = *input;
+    thrust::detail::it_value_t<InputIt> x = *input;
     return op(x);
   }
 
diff --git a/thrust/thrust/system/cuda/detail/for_each.h b/thrust/thrust/system/cuda/detail/for_each.h
index d0a3a65032a..ef84f31c99e 100644
--- a/thrust/thrust/system/cuda/detail/for_each.h
+++ b/thrust/thrust/system/cuda/detail/for_each.h
@@ -72,7 +72,7 @@ Input THRUST_FUNCTION for_each_n(execution_policy<Derived>& policy, Input first,
 template <class Derived, class Input, class UnaryOp>
 Input THRUST_FUNCTION for_each(execution_policy<Derived>& policy, Input first, Input last, UnaryOp op)
 {
-  using size_type = typename iterator_traits<Input>::difference_type;
+  using size_type = thrust::detail::it_difference_t<Input>;
   size_type count = static_cast<size_type>(thrust::distance(first, last));
 
   return THRUST_NS_QUALIFIER::cuda_cub::for_each_n(policy, first, count, op);
diff --git a/thrust/thrust/system/cuda/detail/get_value.h b/thrust/thrust/system/cuda/detail/get_value.h
index f9d896e9456..8e7b46dd005 100644
--- a/thrust/thrust/system/cuda/detail/get_value.h
+++ b/thrust/thrust/system/cuda/detail/get_value.h
@@ -39,7 +39,7 @@ THRUST_NAMESPACE_BEGIN
 namespace cuda_cub
 {
 template <typename DerivedPolicy, typename Pointer>
-_CCCL_HOST_DEVICE iterator_value_t<Pointer> get_value(execution_policy<DerivedPolicy>& exec, Pointer ptr)
+_CCCL_HOST_DEVICE thrust::detail::it_value_t<Pointer> get_value(execution_policy<DerivedPolicy>& exec, Pointer ptr)
 {
   // Because of https://docs.nvidia.com/cuda/cuda-c-programming-guide/#cuda-arch point 2., if a call from a __host__
   // __device__ function leads to the template instantiation of a __global__ function, then this instantiation needs to
@@ -50,7 +50,7 @@ _CCCL_HOST_DEVICE iterator_value_t<Pointer> get_value(execution_policy<DerivedPo
     _CCCL_HOST auto operator()(execution_policy<DerivedPolicy>& exec, Pointer ptr)
     {
       // implemented with assign_value, which requires a type with a default constructor
-      iterator_value_t<Pointer> result;
+      thrust::detail::it_value_t<Pointer> result;
       host_system_tag host_tag;
       cross_system<host_system_tag, DerivedPolicy> systems(host_tag, exec);
       assign_value(systems, &result, ptr);
diff --git a/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h b/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h
index a27e9e48a87..c2249bcc85e 100644
--- a/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h
+++ b/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h
@@ -85,7 +85,7 @@ OutputIt _CCCL_HOST cross_system_copy_n(
   thrust::detail::true_type) // trivial copy
 
 {
-  using InputTy = typename iterator_traits<InputIt>::value_type;
+  using InputTy = thrust::detail::it_value_t<InputIt>;
   if (n > 0)
   {
     trivial_device_copy(
@@ -110,7 +110,7 @@ OutputIt _CCCL_HOST cross_system_copy_n(
   thrust::detail::false_type) // non-trivial copy
 {
   // get type of the input data
-  using InputTy = typename thrust::iterator_value<InputIt>::type;
+  using InputTy = thrust::detail::it_value_t<InputIt>;
 
   // copy input data into host temp storage
   InputIt last = first;
@@ -152,7 +152,7 @@ OutputIt _CCCL_HOST cross_system_copy_n(
 
 {
   // get type of the input data
-  using InputTy = typename thrust::iterator_value<InputIt>::type;
+  using InputTy = thrust::detail::it_value_t<InputIt>;
 
   // allocate device temp storage
   thrust::detail::temporary_array<InputTy, D> d_in_ptr(device_s, num_items);
diff --git a/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h b/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
index 0cb64431956..9b1e1cd7457 100644
--- a/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
+++ b/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
@@ -61,7 +61,7 @@ template <class Derived, class InputIt, class OutputIt>
 OutputIt THRUST_RUNTIME_FUNCTION device_to_device(
   execution_policy<Derived>& policy, InputIt first, InputIt last, OutputIt result, thrust::detail::true_type)
 {
-  using InputTy = typename thrust::iterator_traits<InputIt>::value_type;
+  using InputTy = thrust::detail::it_value_t<InputIt>;
   const auto n  = thrust::distance(first, last);
   if (n > 0)
   {
diff --git a/thrust/thrust/system/cuda/detail/merge.h b/thrust/thrust/system/cuda/detail/merge.h
index 0c39a15f669..c6b480e6c88 100644
--- a/thrust/thrust/system/cuda/detail/merge.h
+++ b/thrust/thrust/system/cuda/detail/merge.h
@@ -64,7 +64,7 @@ merge(execution_policy<Derived>& policy,
       CompareOp compare_op = {})
 {
   THRUST_CDP_DISPATCH(
-    (using size_type         = typename iterator_traits<KeysIt1>::difference_type;
+    (using size_type         = thrust::detail::it_difference_t<KeysIt1>;
      const auto num_keys1    = static_cast<size_type>(thrust::distance(keys1_begin, keys1_end));
      const auto num_keys2    = static_cast<size_type>(thrust::distance(keys2_begin, keys2_end));
      const auto num_keys_out = num_keys1 + num_keys2;
@@ -151,7 +151,7 @@ pair<KeysOutputIt, ItemsOutputIt> _CCCL_HOST_DEVICE merge_by_key(
   CompareOp compare_op = {})
 {
   THRUST_CDP_DISPATCH(
-    (using size_type = typename iterator_traits<KeysIt1>::difference_type;
+    (using size_type = thrust::detail::it_difference_t<KeysIt1>;
 
      const auto num_keys1    = static_cast<size_type>(thrust::distance(keys1_begin, keys1_end));
      const auto num_keys2    = static_cast<size_type>(thrust::distance(keys2_begin, keys2_end));
diff --git a/thrust/thrust/system/cuda/detail/mismatch.h b/thrust/thrust/system/cuda/detail/mismatch.h
index 63e64b722ec..6e723bf4d75 100644
--- a/thrust/thrust/system/cuda/detail/mismatch.h
+++ b/thrust/thrust/system/cuda/detail/mismatch.h
@@ -69,7 +69,7 @@ template <class ValueType, class InputIt1, class InputIt2, class BinaryOp>
 struct transform_pair_of_input_iterators_t
 {
   using self_t            = transform_pair_of_input_iterators_t;
-  using difference_type   = typename iterator_traits<InputIt1>::difference_type;
+  using difference_type   = thrust::detail::it_difference_t<InputIt1>;
   using value_type        = ValueType;
   using pointer           = void;
   using reference         = value_type;
@@ -208,7 +208,7 @@ template <class Derived, class InputIt1, class InputIt2>
 pair<InputIt1, InputIt2> _CCCL_HOST_DEVICE
 mismatch(execution_policy<Derived>& policy, InputIt1 first1, InputIt1 last1, InputIt2 first2)
 {
-  using InputType1 = typename thrust::iterator_value<InputIt1>::type;
+  using InputType1 = thrust::detail::it_value_t<InputIt1>;
   return cuda_cub::mismatch(policy, first1, last1, first2, equal_to<InputType1>());
 }
 
diff --git a/thrust/thrust/system/cuda/detail/partition.h b/thrust/thrust/system/cuda/detail/partition.h
index 690e1e4dd6d..d016dae268b 100644
--- a/thrust/thrust/system/cuda/detail/partition.h
+++ b/thrust/thrust/system/cuda/detail/partition.h
@@ -166,7 +166,7 @@ THRUST_RUNTIME_FUNCTION std::size_t partition(
   OutputIt output,
   Predicate predicate)
 {
-  using size_type = typename iterator_traits<InputIt>::difference_type;
+  using size_type = thrust::detail::it_difference_t<InputIt>;
 
   size_type num_items = thrust::distance(first, last);
   std::size_t num_selected{};
@@ -241,7 +241,7 @@ THRUST_RUNTIME_FUNCTION InputIt inplace_partition(
   }
 
   // Element type of the input iterator
-  using value_t         = typename iterator_traits<InputIt>::value_type;
+  using value_t         = thrust::detail::it_value_t<InputIt>;
   std::size_t num_items = static_cast<std::size_t>(thrust::distance(first, last));
 
   // Allocate temporary storage, which will serve as the input to the partition
diff --git a/thrust/thrust/system/cuda/detail/reduce.h b/thrust/thrust/system/cuda/detail/reduce.h
index 0c266f71b54..e5014106045 100644
--- a/thrust/thrust/system/cuda/detail/reduce.h
+++ b/thrust/thrust/system/cuda/detail/reduce.h
@@ -46,7 +46,6 @@
 #  include <thrust/detail/alignment.h>
 #  include <thrust/detail/raw_reference_cast.h>
 #  include <thrust/detail/temporary_array.h>
-#  include <thrust/detail/type_traits/iterator/is_output_iterator.h>
 #  include <thrust/distance.h>
 #  include <thrust/functional.h>
 #  include <thrust/system/cuda/detail/cdp_dispatch.h>
@@ -857,7 +856,7 @@ reduce_n(execution_policy<Derived>& policy, InputIt first, Size num_items, T ini
 template <class Derived, class InputIt, class T, class BinaryOp>
 _CCCL_HOST_DEVICE T reduce(execution_policy<Derived>& policy, InputIt first, InputIt last, T init, BinaryOp binary_op)
 {
-  using size_type = typename iterator_traits<InputIt>::difference_type;
+  using size_type = thrust::detail::it_difference_t<InputIt>;
   // FIXME: Check for RA iterator.
   size_type num_items = static_cast<size_type>(thrust::distance(first, last));
   return cuda_cub::reduce_n(policy, first, num_items, init, binary_op);
@@ -870,10 +869,10 @@ _CCCL_HOST_DEVICE T reduce(execution_policy<Derived>& policy, InputIt first, Inp
 }
 
 template <class Derived, class InputIt>
-_CCCL_HOST_DEVICE typename iterator_traits<InputIt>::value_type
+_CCCL_HOST_DEVICE thrust::detail::it_value_t<InputIt>
 reduce(execution_policy<Derived>& policy, InputIt first, InputIt last)
 {
-  using value_type = typename iterator_traits<InputIt>::value_type;
+  using value_type = thrust::detail::it_value_t<InputIt>;
   return cuda_cub::reduce(policy, first, last, value_type(0));
 }
 
diff --git a/thrust/thrust/system/cuda/detail/reduce_by_key.h b/thrust/thrust/system/cuda/detail/reduce_by_key.h
index 902d5574d91..17f01006851 100644
--- a/thrust/thrust/system/cuda/detail/reduce_by_key.h
+++ b/thrust/thrust/system/cuda/detail/reduce_by_key.h
@@ -58,6 +58,8 @@
 #  include <thrust/system/cuda/detail/par_to_seq.h>
 #  include <thrust/system/cuda/detail/util.h>
 
+#  include <cuda/std/iterator>
+
 #  include <cstdint>
 
 THRUST_NAMESPACE_BEGIN
@@ -148,8 +150,8 @@ template <class KeysInputIt,
           class Size>
 struct ReduceByKeyAgent
 {
-  using key_type   = typename iterator_traits<KeysInputIt>::value_type;
-  using value_type = typename iterator_traits<ValuesInputIt>::value_type;
+  using key_type   = thrust::detail::it_value_t<KeysInputIt>;
+  using value_type = thrust::detail::it_value_t<ValuesInputIt>;
   using size_type  = Size;
 
   using size_value_pair_t = cub::KeyValuePair<size_type, value_type>;
@@ -891,7 +893,7 @@ THRUST_RUNTIME_FUNCTION pair<KeysOutputIt, ValuesOutputIt> reduce_by_key(
   EqualityOp equality_op,
   ReductionOp reduction_op)
 {
-  using size_type = typename iterator_traits<KeysInputIt>::difference_type;
+  using size_type = thrust::detail::it_difference_t<KeysInputIt>;
 
   size_type num_items = thrust::distance(keys_first, keys_last);
 
@@ -961,9 +963,9 @@ pair<KeyOutputIt, ValOutputIt> _CCCL_HOST_DEVICE reduce_by_key(
   ValOutputIt values_output,
   BinaryPred binary_pred)
 {
-  using value_type = typename thrust::detail::eval_if<thrust::detail::is_output_iterator<ValOutputIt>::value,
-                                                      thrust::iterator_value<ValInputIt>,
-                                                      thrust::iterator_value<ValOutputIt>>::type;
+  using value_type = ::cuda::std::_If<thrust::detail::is_output_iterator<ValOutputIt>,
+                                      thrust::detail::it_value_t<ValInputIt>,
+                                      thrust::detail::it_value_t<ValOutputIt>>;
   return cuda_cub::reduce_by_key(
     policy, keys_first, keys_last, values_first, keys_output, values_output, binary_pred, plus<value_type>());
 }
@@ -977,7 +979,7 @@ pair<KeyOutputIt, ValOutputIt> _CCCL_HOST_DEVICE reduce_by_key(
   KeyOutputIt keys_output,
   ValOutputIt values_output)
 {
-  using KeyT = typename thrust::iterator_value<KeyInputIt>::type;
+  using KeyT = thrust::detail::it_value_t<KeyInputIt>;
   return cuda_cub::reduce_by_key(
     policy, keys_first, keys_last, values_first, keys_output, values_output, equal_to<KeyT>());
 }
diff --git a/thrust/thrust/system/cuda/detail/replace.h b/thrust/thrust/system/cuda/detail/replace.h
index 45682ac52c2..12257f2f224 100644
--- a/thrust/thrust/system/cuda/detail/replace.h
+++ b/thrust/thrust/system/cuda/detail/replace.h
@@ -127,7 +127,7 @@ OutputIt _CCCL_HOST_DEVICE replace_copy_if(
   Predicate predicate,
   T const& new_value)
 {
-  using output_type    = typename iterator_traits<OutputIt>::value_type;
+  using output_type    = thrust::detail::it_value_t<OutputIt>;
   using new_value_if_t = __replace::new_value_if_f<Predicate, T, output_type>;
   return cuda_cub::transform(policy, first, last, result, new_value_if_t(predicate, new_value));
 }
@@ -142,7 +142,7 @@ OutputIt _CCCL_HOST_DEVICE replace_copy_if(
   Predicate predicate,
   T const& new_value)
 {
-  using output_type    = typename iterator_traits<OutputIt>::value_type;
+  using output_type    = thrust::detail::it_value_t<OutputIt>;
   using new_value_if_t = __replace::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/cuda/detail/reverse.h b/thrust/thrust/system/cuda/detail/reverse.h
index 57cc16852fd..682b0efcf5f 100644
--- a/thrust/thrust/system/cuda/detail/reverse.h
+++ b/thrust/thrust/system/cuda/detail/reverse.h
@@ -58,6 +58,8 @@ THRUST_NAMESPACE_END
 #  include <thrust/system/cuda/detail/copy.h>
 #  include <thrust/system/cuda/detail/swap_ranges.h>
 
+#  include <cuda/std/iterator>
+
 THRUST_NAMESPACE_BEGIN
 namespace cuda_cub
 {
@@ -71,7 +73,7 @@ ResultIt _CCCL_HOST_DEVICE reverse_copy(execution_policy<Derived>& policy, Items
 template <class Derived, class ItemsIt>
 void _CCCL_HOST_DEVICE reverse(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last)
 {
-  using difference_type = typename thrust::iterator_difference<ItemsIt>::type;
+  using difference_type = thrust::detail::it_difference_t<ItemsIt>;
 
   // find the midpoint of [first,last)
   difference_type N = thrust::distance(first, last);
diff --git a/thrust/thrust/system/cuda/detail/scan.h b/thrust/thrust/system/cuda/detail/scan.h
index 5402437673f..6f83aafcf21 100644
--- a/thrust/thrust/system/cuda/detail/scan.h
+++ b/thrust/thrust/system/cuda/detail/scan.h
@@ -63,7 +63,7 @@ template <typename Derived, typename InputIt, typename Size, typename OutputIt,
 _CCCL_HOST_DEVICE OutputIt inclusive_scan_n_impl(
   thrust::cuda_cub::execution_policy<Derived>& policy, InputIt first, Size num_items, OutputIt result, ScanOp scan_op)
 {
-  using AccumT     = typename thrust::iterator_traits<InputIt>::value_type;
+  using AccumT     = thrust::detail::it_value_t<InputIt>;
   using Dispatch32 = cub::DispatchScan<InputIt, OutputIt, ScanOp, cub::NullType, std::uint32_t, AccumT>;
   using Dispatch64 = cub::DispatchScan<InputIt, OutputIt, ScanOp, cub::NullType, std::uint64_t, AccumT>;
 
@@ -262,7 +262,7 @@ template <typename Derived, typename InputIt, typename OutputIt, typename ScanOp
 _CCCL_HOST_DEVICE OutputIt inclusive_scan(
   thrust::cuda_cub::execution_policy<Derived>& policy, InputIt first, InputIt last, OutputIt result, ScanOp scan_op)
 {
-  using diff_t           = typename thrust::iterator_traits<InputIt>::difference_type;
+  using diff_t           = thrust::detail::it_difference_t<InputIt>;
   diff_t const num_items = thrust::distance(first, last);
   return thrust::cuda_cub::inclusive_scan_n(policy, first, num_items, result, scan_op);
 }
@@ -276,7 +276,7 @@ _CCCL_HOST_DEVICE OutputIt inclusive_scan(
   T init,
   ScanOp scan_op)
 {
-  using diff_t           = typename thrust::iterator_traits<InputIt>::difference_type;
+  using diff_t           = thrust::detail::it_difference_t<InputIt>;
   diff_t const num_items = thrust::distance(first, last);
   return thrust::cuda_cub::inclusive_scan_n(policy, first, num_items, result, init, scan_op);
 }
@@ -314,7 +314,7 @@ _CCCL_HOST_DEVICE OutputIt exclusive_scan(
   T init,
   ScanOp scan_op)
 {
-  using diff_t           = typename thrust::iterator_traits<InputIt>::difference_type;
+  using diff_t           = thrust::detail::it_difference_t<InputIt>;
   diff_t const num_items = thrust::distance(first, last);
   return thrust::cuda_cub::exclusive_scan_n(policy, first, num_items, result, init, scan_op);
 }
@@ -330,7 +330,7 @@ template <typename Derived, typename InputIt, typename OutputIt>
 _CCCL_HOST_DEVICE OutputIt
 exclusive_scan(thrust::cuda_cub::execution_policy<Derived>& policy, InputIt first, InputIt last, OutputIt result)
 {
-  using init_type = typename thrust::iterator_traits<InputIt>::value_type;
+  using init_type = thrust::detail::it_value_t<InputIt>;
   return cuda_cub::exclusive_scan(policy, first, last, result, init_type{});
 };
 
diff --git a/thrust/thrust/system/cuda/detail/scan_by_key.h b/thrust/thrust/system/cuda/detail/scan_by_key.h
index 313f69d7758..2edb0ba7e41 100644
--- a/thrust/thrust/system/cuda/detail/scan_by_key.h
+++ b/thrust/thrust/system/cuda/detail/scan_by_key.h
@@ -89,7 +89,7 @@ _CCCL_HOST_DEVICE ValuesOutIt inclusive_scan_by_key_n(
   using KeysInUnwrapIt    = thrust::try_unwrap_contiguous_iterator_t<KeysInIt>;
   using ValuesInUnwrapIt  = thrust::try_unwrap_contiguous_iterator_t<ValuesInIt>;
   using ValuesOutUnwrapIt = thrust::try_unwrap_contiguous_iterator_t<ValuesOutIt>;
-  using AccumT            = typename thrust::iterator_traits<ValuesInUnwrapIt>::value_type;
+  using AccumT            = thrust::detail::it_value_t<ValuesInUnwrapIt>;
 
   auto keys_unwrap   = thrust::try_unwrap_contiguous_iterator(keys);
   auto values_unwrap = thrust::try_unwrap_contiguous_iterator(values);
@@ -395,7 +395,7 @@ ValOutputIt _CCCL_HOST_DEVICE exclusive_scan_by_key(
   ValInputIt value_first,
   ValOutputIt value_result)
 {
-  using value_type = typename thrust::iterator_traits<ValInputIt>::value_type;
+  using value_type = thrust::detail::it_value_t<ValInputIt>;
   return cuda_cub::exclusive_scan_by_key(policy, key_first, key_last, value_first, value_result, value_type{});
 }
 
diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h
index 85d03ae51cb..28c54146190 100644
--- a/thrust/thrust/system/cuda/detail/set_operations.h
+++ b/thrust/thrust/system/cuda/detail/set_operations.h
@@ -128,7 +128,7 @@ THRUST_DEVICE_FUNCTION Size biased_binary_search(It data, Size count, T key, Int
 template <bool UpperBound, class Size, class It1, class It2, class Comp>
 THRUST_DEVICE_FUNCTION Size merge_path(It1 a, Size aCount, It2 b, Size bCount, Size diag, Comp comp)
 {
-  using T = typename thrust::iterator_traits<It1>::value_type;
+  using T = thrust::detail::it_value_t<It1>;
 
   Size begin = ::cuda::std::max<Size>(0, diag - bCount);
   Size end   = ::cuda::std::min<Size>(diag, aCount);
@@ -155,7 +155,7 @@ template <class It1, class It2, class Size, class Size2, class CompareOp>
 THRUST_DEVICE_FUNCTION pair<Size, Size>
 balanced_path(It1 keys1, It2 keys2, Size num_keys1, Size num_keys2, Size diag, Size2 levels, CompareOp compare_op)
 {
-  using T = typename iterator_traits<It1>::value_type;
+  using T = thrust::detail::it_value_t<It1>;
 
   Size index1 = merge_path<false>(keys1, num_keys1, keys2, num_keys2, diag, compare_op);
   Size index2 = diag - index1;
@@ -275,10 +275,10 @@ template <class KeysIt1,
           class HAS_VALUES>
 struct SetOpAgent
 {
-  using key1_type   = typename iterator_traits<KeysIt1>::value_type;
-  using key2_type   = typename iterator_traits<KeysIt2>::value_type;
-  using value1_type = typename iterator_traits<ValuesIt1>::value_type;
-  using value2_type = typename iterator_traits<ValuesIt2>::value_type;
+  using key1_type   = thrust::detail::it_value_t<KeysIt1>;
+  using key2_type   = thrust::detail::it_value_t<KeysIt2>;
+  using value1_type = thrust::detail::it_value_t<ValuesIt1>;
+  using value2_type = thrust::detail::it_value_t<ValuesIt2>;
 
   using key_type   = key1_type;
   using value_type = value1_type;
@@ -1149,7 +1149,7 @@ THRUST_RUNTIME_FUNCTION pair<KeysOutputIt, ValuesOutputIt> set_operations(
   CompareOp compare_op,
   SetOp set_op)
 {
-  using size_type = typename iterator_traits<KeysIt1>::difference_type;
+  using size_type = thrust::detail::it_difference_t<KeysIt1>;
 
   size_type num_keys1 = static_cast<size_type>(thrust::distance(keys1_first, keys1_last));
   size_type num_keys2 = static_cast<size_type>(thrust::distance(keys2_first, keys2_last));
@@ -1247,7 +1247,7 @@ OutputIt _CCCL_HOST_DEVICE set_difference(
   CompareOp compare)
 {
   THRUST_CDP_DISPATCH(
-    (using items1_t = thrust::iterator_value_t<ItemsIt1>; items1_t* null_ = nullptr;
+    (using items1_t = thrust::detail::it_value_t<ItemsIt1>; items1_t* null_ = nullptr;
      auto tmp = __set_operations::set_operations<thrust::detail::false_type>(
        policy,
        items1_first,
@@ -1275,7 +1275,7 @@ OutputIt _CCCL_HOST_DEVICE set_difference(
   ItemsIt2 items2_last,
   OutputIt result)
 {
-  using value_type = typename thrust::iterator_value<ItemsIt1>::type;
+  using value_type = thrust::detail::it_value_t<ItemsIt1>;
   return cuda_cub::set_difference(
     policy, items1_first, items1_last, items2_first, items2_last, result, less<value_type>());
 }
@@ -1294,7 +1294,7 @@ OutputIt _CCCL_HOST_DEVICE set_intersection(
   CompareOp compare)
 {
   THRUST_CDP_DISPATCH(
-    (using items1_t = thrust::iterator_value_t<ItemsIt1>; items1_t* null_ = nullptr;
+    (using items1_t = thrust::detail::it_value_t<ItemsIt1>; items1_t* null_ = nullptr;
      auto tmp = __set_operations::set_operations<thrust::detail::false_type>(
        policy,
        items1_first,
@@ -1322,7 +1322,7 @@ OutputIt _CCCL_HOST_DEVICE set_intersection(
   ItemsIt2 items2_last,
   OutputIt result)
 {
-  using value_type = typename thrust::iterator_value<ItemsIt1>::type;
+  using value_type = thrust::detail::it_value_t<ItemsIt1>;
   return cuda_cub::set_intersection(
     policy, items1_first, items1_last, items2_first, items2_last, result, less<value_type>());
 }
@@ -1341,7 +1341,7 @@ OutputIt _CCCL_HOST_DEVICE set_symmetric_difference(
   CompareOp compare)
 {
   THRUST_CDP_DISPATCH(
-    (using items1_t = thrust::iterator_value_t<ItemsIt1>; items1_t* null_ = nullptr;
+    (using items1_t = thrust::detail::it_value_t<ItemsIt1>; items1_t* null_ = nullptr;
      auto tmp = __set_operations::set_operations<thrust::detail::false_type>(
        policy,
        items1_first,
@@ -1369,7 +1369,7 @@ OutputIt _CCCL_HOST_DEVICE set_symmetric_difference(
   ItemsIt2 items2_last,
   OutputIt result)
 {
-  using value_type = typename thrust::iterator_value<ItemsIt1>::type;
+  using value_type = thrust::detail::it_value_t<ItemsIt1>;
   return cuda_cub::set_symmetric_difference(
     policy, items1_first, items1_last, items2_first, items2_last, result, less<value_type>());
 }
@@ -1388,7 +1388,7 @@ OutputIt _CCCL_HOST_DEVICE set_union(
   CompareOp compare)
 {
   THRUST_CDP_DISPATCH(
-    (using items1_t = thrust::iterator_value_t<ItemsIt1>; items1_t* null_ = nullptr;
+    (using items1_t = thrust::detail::it_value_t<ItemsIt1>; items1_t* null_ = nullptr;
      auto tmp = __set_operations::set_operations<thrust::detail::false_type>(
        policy,
        items1_first,
@@ -1416,7 +1416,7 @@ OutputIt _CCCL_HOST_DEVICE set_union(
   ItemsIt2 items2_last,
   OutputIt result)
 {
-  using value_type = typename thrust::iterator_value<ItemsIt1>::type;
+  using value_type = thrust::detail::it_value_t<ItemsIt1>;
   return cuda_cub::set_union(policy, items1_first, items1_last, items2_first, items2_last, result, less<value_type>());
 }
 
@@ -1489,7 +1489,7 @@ pair<KeysOutputIt, ItemsOutputIt> _CCCL_HOST_DEVICE set_difference_by_key(
   KeysOutputIt keys_result,
   ItemsOutputIt items_result)
 {
-  using value_type = typename thrust::iterator_value<KeysIt1>::type;
+  using value_type = thrust::detail::it_value_t<KeysIt1>;
   return cuda_cub::set_difference_by_key(
     policy,
     keys1_first,
@@ -1563,7 +1563,7 @@ pair<KeysOutputIt, ItemsOutputIt> _CCCL_HOST_DEVICE set_intersection_by_key(
   KeysOutputIt keys_result,
   ItemsOutputIt items_result)
 {
-  using value_type = typename thrust::iterator_value<KeysIt1>::type;
+  using value_type = thrust::detail::it_value_t<KeysIt1>;
   return cuda_cub::set_intersection_by_key(
     policy,
     keys1_first,
@@ -1639,7 +1639,7 @@ pair<KeysOutputIt, ItemsOutputIt> _CCCL_HOST_DEVICE set_symmetric_difference_by_
   KeysOutputIt keys_result,
   ItemsOutputIt items_result)
 {
-  using value_type = typename thrust::iterator_value<KeysIt1>::type;
+  using value_type = thrust::detail::it_value_t<KeysIt1>;
   return cuda_cub::set_symmetric_difference_by_key(
     policy,
     keys1_first,
@@ -1716,7 +1716,7 @@ pair<KeysOutputIt, ItemsOutputIt> _CCCL_HOST_DEVICE set_union_by_key(
   KeysOutputIt keys_result,
   ItemsOutputIt items_result)
 {
-  using value_type = typename thrust::iterator_value<KeysIt1>::type;
+  using value_type = thrust::detail::it_value_t<KeysIt1>;
   return cuda_cub::set_union_by_key(
     policy,
     keys1_first,
diff --git a/thrust/thrust/system/cuda/detail/sort.h b/thrust/thrust/system/cuda/detail/sort.h
index 6618e5fc739..a32e5bd0b75 100644
--- a/thrust/thrust/system/cuda/detail/sort.h
+++ b/thrust/thrust/system/cuda/detail/sort.h
@@ -160,7 +160,7 @@ THRUST_RUNTIME_FUNCTION void merge_sort(
   execution_policy<Derived>& policy, KeysIt keys_first, KeysIt keys_last, ItemsIt items_first, CompareOp compare_op)
 
 {
-  using size_type = typename iterator_traits<KeysIt>::difference_type;
+  using size_type = thrust::detail::it_difference_t<KeysIt>;
 
   size_type count = static_cast<size_type>(thrust::distance(keys_first, keys_last));
 
@@ -347,7 +347,7 @@ template <
   class KeysIt,
   class ItemsIt,
   class CompareOp,
-  ::cuda::std::enable_if_t<!can_use_primitive_sort<typename iterator_value<KeysIt>::type, CompareOp>::value, int> = 0>
+  ::cuda::std::enable_if_t<!can_use_primitive_sort<thrust::detail::it_value_t<KeysIt>, CompareOp>::value, int> = 0>
 THRUST_RUNTIME_FUNCTION void
 smart_sort(Policy& policy, KeysIt keys_first, KeysIt keys_last, ItemsIt items_first, CompareOp compare_op)
 {
@@ -361,7 +361,7 @@ template <
   class KeysIt,
   class ItemsIt,
   class CompareOp,
-  ::cuda::std::enable_if_t<can_use_primitive_sort<typename iterator_value<KeysIt>::type, CompareOp>::value, int> = 0>
+  ::cuda::std::enable_if_t<can_use_primitive_sort<thrust::detail::it_value_t<KeysIt>, CompareOp>::value, int> = 0>
 THRUST_RUNTIME_FUNCTION void smart_sort(
   execution_policy<Policy>& policy, KeysIt keys_first, KeysIt keys_last, ItemsIt items_first, CompareOp compare_op)
 {
@@ -413,7 +413,7 @@ _CCCL_EXEC_CHECK_DISABLE
 template <class Derived, class ItemsIt, class CompareOp>
 void _CCCL_HOST_DEVICE sort(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, CompareOp compare_op)
 {
-  THRUST_CDP_DISPATCH((using item_t = thrust::iterator_value_t<ItemsIt>; item_t* null_ = nullptr;
+  THRUST_CDP_DISPATCH((using item_t = thrust::detail::it_value_t<ItemsIt>; item_t* null_ = nullptr;
                        __smart_sort::smart_sort<thrust::detail::false_type, thrust::detail::false_type>(
                          policy, first, last, null_, compare_op);),
                       (thrust::sort(cvt_to_seq(derived_cast(policy)), first, last, compare_op);));
@@ -423,7 +423,7 @@ _CCCL_EXEC_CHECK_DISABLE
 template <class Derived, class ItemsIt, class CompareOp>
 void _CCCL_HOST_DEVICE stable_sort(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last, CompareOp compare_op)
 {
-  THRUST_CDP_DISPATCH((using item_t = thrust::iterator_value_t<ItemsIt>; item_t* null_ = nullptr;
+  THRUST_CDP_DISPATCH((using item_t = thrust::detail::it_value_t<ItemsIt>; item_t* null_ = nullptr;
                        __smart_sort::smart_sort<thrust::detail::false_type, thrust::detail::true_type>(
                          policy, first, last, null_, compare_op);),
                       (thrust::stable_sort(cvt_to_seq(derived_cast(policy)), first, last, compare_op);));
@@ -456,14 +456,14 @@ void _CCCL_HOST_DEVICE stable_sort_by_key(
 template <class Derived, class ItemsIt>
 void _CCCL_HOST_DEVICE sort(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last)
 {
-  using item_type = typename thrust::iterator_value<ItemsIt>::type;
+  using item_type = thrust::detail::it_value_t<ItemsIt>;
   cuda_cub::sort(policy, first, last, less<item_type>());
 }
 
 template <class Derived, class ItemsIt>
 void _CCCL_HOST_DEVICE stable_sort(execution_policy<Derived>& policy, ItemsIt first, ItemsIt last)
 {
-  using item_type = typename thrust::iterator_value<ItemsIt>::type;
+  using item_type = thrust::detail::it_value_t<ItemsIt>;
   cuda_cub::stable_sort(policy, first, last, less<item_type>());
 }
 
@@ -471,7 +471,7 @@ template <class Derived, class KeysIt, class ValuesIt>
 void _CCCL_HOST_DEVICE
 sort_by_key(execution_policy<Derived>& policy, KeysIt keys_first, KeysIt keys_last, ValuesIt values)
 {
-  using key_type = typename thrust::iterator_value<KeysIt>::type;
+  using key_type = thrust::detail::it_value_t<KeysIt>;
   cuda_cub::sort_by_key(policy, keys_first, keys_last, values, less<key_type>());
 }
 
@@ -479,7 +479,7 @@ template <class Derived, class KeysIt, class ValuesIt>
 void _CCCL_HOST_DEVICE
 stable_sort_by_key(execution_policy<Derived>& policy, KeysIt keys_first, KeysIt keys_last, ValuesIt values)
 {
-  using key_type = typename thrust::iterator_value<KeysIt>::type;
+  using key_type = thrust::detail::it_value_t<KeysIt>;
   cuda_cub::stable_sort_by_key(policy, keys_first, keys_last, values, less<key_type>());
 }
 
diff --git a/thrust/thrust/system/cuda/detail/swap_ranges.h b/thrust/thrust/system/cuda/detail/swap_ranges.h
index 73a73d63775..23cd499ad2a 100644
--- a/thrust/thrust/system/cuda/detail/swap_ranges.h
+++ b/thrust/thrust/system/cuda/detail/swap_ranges.h
@@ -61,8 +61,8 @@ struct swap_f
   ItemsIt1 items1;
   ItemsIt2 items2;
 
-  using value1_type = typename iterator_traits<ItemsIt1>::value_type;
-  using value2_type = typename iterator_traits<ItemsIt2>::value_type;
+  using value1_type = thrust::detail::it_value_t<ItemsIt1>;
+  using value2_type = thrust::detail::it_value_t<ItemsIt2>;
 
   THRUST_FUNCTION
   swap_f(ItemsIt1 items1_, ItemsIt2 items2_)
@@ -88,7 +88,7 @@ template <class Derived, class ItemsIt1, class ItemsIt2>
 ItemsIt2 _CCCL_HOST_DEVICE
 swap_ranges(execution_policy<Derived>& policy, ItemsIt1 first1, ItemsIt1 last1, ItemsIt2 first2)
 {
-  using size_type = typename iterator_traits<ItemsIt1>::difference_type;
+  using size_type = thrust::detail::it_difference_t<ItemsIt1>;
 
   size_type num_items = static_cast<size_type>(thrust::distance(first1, last1));
 
diff --git a/thrust/thrust/system/cuda/detail/tabulate.h b/thrust/thrust/system/cuda/detail/tabulate.h
index 6caec210120..dee20887a77 100644
--- a/thrust/thrust/system/cuda/detail/tabulate.h
+++ b/thrust/thrust/system/cuda/detail/tabulate.h
@@ -72,7 +72,7 @@ struct functor
 template <class Derived, class Iterator, class TabulateOp>
 void _CCCL_HOST_DEVICE tabulate(execution_policy<Derived>& policy, Iterator first, Iterator last, TabulateOp tabulate_op)
 {
-  using size_type = typename iterator_traits<Iterator>::difference_type;
+  using size_type = thrust::detail::it_difference_t<Iterator>;
 
   size_type count = thrust::distance(first, last);
 
diff --git a/thrust/thrust/system/cuda/detail/transform.h b/thrust/thrust/system/cuda/detail/transform.h
index 2baf4810550..6863fba4490 100644
--- a/thrust/thrust/system/cuda/detail/transform.h
+++ b/thrust/thrust/system/cuda/detail/transform.h
@@ -319,7 +319,7 @@ OutputIt THRUST_FUNCTION transform_if(
   TransformOp transform_op,
   Predicate predicate)
 {
-  using size_type     = typename iterator_traits<InputIt>::difference_type;
+  using size_type     = thrust::detail::it_difference_t<InputIt>;
   size_type num_items = static_cast<size_type>(thrust::distance(first, last));
   return __transform::unary(policy, first, result, num_items, stencil, transform_op, predicate);
 } // func transform_if
@@ -341,7 +341,7 @@ OutputIt THRUST_FUNCTION
 transform(execution_policy<Derived>& policy, InputIt first, InputIt last, OutputIt result, TransformOp transform_op)
 {
   THRUST_CDP_DISPATCH(
-    (using size_type      = typename iterator_traits<InputIt>::difference_type;
+    (using size_type      = thrust::detail::it_difference_t<InputIt>;
      const auto num_items = static_cast<size_type>(thrust::distance(first, last));
      return __transform::cub_transform_many(policy, ::cuda::std::make_tuple(first), result, num_items, transform_op);),
     (while (first != last) {
@@ -372,7 +372,7 @@ OutputIt THRUST_FUNCTION transform_if(
   TransformOp transform_op,
   Predicate predicate)
 {
-  using size_type     = typename iterator_traits<InputIt1>::difference_type;
+  using size_type     = thrust::detail::it_difference_t<InputIt1>;
   size_type num_items = static_cast<size_type>(thrust::distance(first1, last1));
   return __transform::binary(policy, first1, first2, result, num_items, stencil, transform_op, predicate);
 } // func transform_if
@@ -387,7 +387,7 @@ OutputIt THRUST_FUNCTION transform(
   TransformOp transform_op)
 {
   THRUST_CDP_DISPATCH(
-    (using size_type      = typename iterator_traits<InputIt1>::difference_type;
+    (using size_type      = thrust::detail::it_difference_t<InputIt1>;
      const auto num_items = static_cast<size_type>(thrust::distance(first1, last1));
      return __transform::cub_transform_many(
        policy, ::cuda::std::make_tuple(first1, first2), result, num_items, transform_op);),
diff --git a/thrust/thrust/system/cuda/detail/transform_reduce.h b/thrust/thrust/system/cuda/detail/transform_reduce.h
index fbb3054b0bf..ed7d26942f7 100644
--- a/thrust/thrust/system/cuda/detail/transform_reduce.h
+++ b/thrust/thrust/system/cuda/detail/transform_reduce.h
@@ -45,7 +45,6 @@
 #  include <thrust/detail/alignment.h>
 #  include <thrust/detail/raw_reference_cast.h>
 #  include <thrust/detail/temporary_array.h>
-#  include <thrust/detail/type_traits/iterator/is_output_iterator.h>
 #  include <thrust/distance.h>
 #  include <thrust/functional.h>
 #  include <thrust/system/cuda/detail/cdp_dispatch.h>
@@ -130,7 +129,7 @@ template <class Derived, class InputIt, class TransformOp, class T, class Reduce
 T _CCCL_HOST_DEVICE transform_reduce(
   execution_policy<Derived>& policy, InputIt first, InputIt last, TransformOp transform_op, T init, ReduceOp reduce_op)
 {
-  using size_type           = typename iterator_traits<InputIt>::difference_type;
+  using size_type           = thrust::detail::it_difference_t<InputIt>;
   const size_type num_items = static_cast<size_type>(thrust::distance(first, last));
 
   THRUST_CDP_DISPATCH(
diff --git a/thrust/thrust/system/cuda/detail/transform_scan.h b/thrust/thrust/system/cuda/detail/transform_scan.h
index 2227249ab16..335e4ea97c3 100644
--- a/thrust/thrust/system/cuda/detail/transform_scan.h
+++ b/thrust/thrust/system/cuda/detail/transform_scan.h
@@ -61,11 +61,11 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan(
   ScanOp scan_op)
 {
   // Use the transformed input iterator's value type per https://wg21.link/P0571
-  using input_type  = typename thrust::iterator_value<InputIt>::type;
+  using input_type  = thrust::detail::it_value_t<InputIt>;
   using result_type = thrust::detail::invoke_result_t<TransformOp, input_type>;
   using value_type  = ::cuda::std::remove_cvref_t<result_type>;
 
-  using size_type              = typename iterator_traits<InputIt>::difference_type;
+  using size_type              = thrust::detail::it_difference_t<InputIt>;
   size_type num_items          = static_cast<size_type>(thrust::distance(first, last));
   using transformed_iterator_t = transform_iterator<TransformOp, InputIt, value_type, value_type>;
 
@@ -82,11 +82,11 @@ OutputIt _CCCL_HOST_DEVICE transform_inclusive_scan(
   InitialValueType init,
   ScanOp scan_op)
 {
-  using input_type  = typename thrust::iterator_value<InputIt>::type;
+  using input_type  = thrust::detail::it_value_t<InputIt>;
   using result_type = thrust::detail::invoke_result_t<TransformOp, input_type>;
   using value_type  = ::cuda::std::remove_cvref_t<result_type>;
 
-  using size_type              = typename iterator_traits<InputIt>::difference_type;
+  using size_type              = thrust::detail::it_difference_t<InputIt>;
   size_type num_items          = static_cast<size_type>(thrust::distance(first, last));
   using transformed_iterator_t = transform_iterator<TransformOp, InputIt, value_type, value_type>;
 
@@ -107,7 +107,7 @@ OutputIt _CCCL_HOST_DEVICE transform_exclusive_scan(
   // Use the initial value type per https://wg21.link/P0571
   using result_type = ::cuda::std::remove_cvref_t<InitialValueType>;
 
-  using size_type              = typename iterator_traits<InputIt>::difference_type;
+  using size_type              = thrust::detail::it_difference_t<InputIt>;
   size_type num_items          = static_cast<size_type>(thrust::distance(first, last));
   using transformed_iterator_t = transform_iterator<TransformOp, InputIt, result_type, result_type>;
 
diff --git a/thrust/thrust/system/cuda/detail/uninitialized_copy.h b/thrust/thrust/system/cuda/detail/uninitialized_copy.h
index d24cf87e6ed..5be6813ebf7 100644
--- a/thrust/thrust/system/cuda/detail/uninitialized_copy.h
+++ b/thrust/thrust/system/cuda/detail/uninitialized_copy.h
@@ -58,8 +58,8 @@ struct functor
   InputIt input;
   OutputIt output;
 
-  using InputType  = typename iterator_traits<InputIt>::value_type;
-  using OutputType = typename iterator_traits<OutputIt>::value_type;
+  using InputType  = thrust::detail::it_value_t<InputIt>;
+  using OutputType = thrust::detail::it_value_t<OutputIt>;
 
   THRUST_FUNCTION
   functor(InputIt input_, OutputIt output_)
diff --git a/thrust/thrust/system/cuda/detail/uninitialized_fill.h b/thrust/thrust/system/cuda/detail/uninitialized_fill.h
index 233f600e2bb..601e0b9f0b1 100644
--- a/thrust/thrust/system/cuda/detail/uninitialized_fill.h
+++ b/thrust/thrust/system/cuda/detail/uninitialized_fill.h
@@ -58,7 +58,7 @@ struct functor
   Iterator items;
   T value;
 
-  using value_type = typename iterator_traits<Iterator>::value_type;
+  using value_type = thrust::detail::it_value_t<Iterator>;
 
   THRUST_FUNCTION
   functor(Iterator items_, T const& value_)
diff --git a/thrust/thrust/system/cuda/detail/unique.h b/thrust/thrust/system/cuda/detail/unique.h
index 7818aa7f916..0d95ffda5fb 100644
--- a/thrust/thrust/system/cuda/detail/unique.h
+++ b/thrust/thrust/system/cuda/detail/unique.h
@@ -73,7 +73,7 @@ _CCCL_HOST_DEVICE OutputIterator unique_copy(
   BinaryPredicate binary_pred);
 
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator> unique_count(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
   ForwardIterator first,
   ForwardIterator last,
@@ -140,7 +140,7 @@ struct Tuning<core::detail::sm52, T>
 template <class ItemsIt, class ItemsOutputIt, class BinaryPred, class Size, class NumSelectedOutIt>
 struct UniqueAgent
 {
-  using item_type = typename iterator_traits<ItemsIt>::value_type;
+  using item_type = thrust::detail::it_value_t<ItemsIt>;
 
   using ScanTileState = cub::ScanTileState<Size>;
 
@@ -524,7 +524,7 @@ THRUST_RUNTIME_FUNCTION ItemsOutputIt unique(
   ItemsOutputIt items_result,
   BinaryPred binary_pred)
 {
-  //  using size_type = typename iterator_traits<ItemsInputIt>::difference_type;
+  //  using size_type = thrust::detail::it_difference_t<ItemsInputIt>;
   using size_type = int;
 
   size_type num_items       = static_cast<size_type>(thrust::distance(items_first, items_last));
@@ -590,7 +590,7 @@ unique_copy(execution_policy<Derived>& policy, InputIt first, InputIt last, Outp
 template <class Derived, class InputIt, class OutputIt>
 OutputIt _CCCL_HOST_DEVICE unique_copy(execution_policy<Derived>& policy, InputIt first, InputIt last, OutputIt result)
 {
-  using input_type = typename iterator_traits<InputIt>::value_type;
+  using input_type = thrust::detail::it_value_t<InputIt>;
   return cuda_cub::unique_copy(policy, first, last, result, equal_to<input_type>());
 }
 
@@ -608,7 +608,7 @@ unique(execution_policy<Derived>& policy, ForwardIt first, ForwardIt last, Binar
 template <class Derived, class ForwardIt>
 ForwardIt _CCCL_HOST_DEVICE unique(execution_policy<Derived>& policy, ForwardIt first, ForwardIt last)
 {
-  using input_type = typename iterator_traits<ForwardIt>::value_type;
+  using input_type = thrust::detail::it_value_t<ForwardIt>;
   return cuda_cub::unique(policy, first, last, equal_to<input_type>());
 }
 
@@ -626,7 +626,7 @@ struct zip_adj_not_predicate
 
 _CCCL_EXEC_CHECK_DISABLE
 template <class Derived, class ForwardIt, class BinaryPred>
-typename thrust::iterator_traits<ForwardIt>::difference_type _CCCL_HOST_DEVICE
+thrust::detail::it_difference_t<ForwardIt> _CCCL_HOST_DEVICE
 unique_count(execution_policy<Derived>& policy, ForwardIt first, ForwardIt last, BinaryPred binary_pred)
 {
   if (first == last)
diff --git a/thrust/thrust/system/cuda/detail/unique_by_key.h b/thrust/thrust/system/cuda/detail/unique_by_key.h
index 468c1bdae6b..fa76ff9ca77 100644
--- a/thrust/thrust/system/cuda/detail/unique_by_key.h
+++ b/thrust/thrust/system/cuda/detail/unique_by_key.h
@@ -182,7 +182,7 @@ THRUST_RUNTIME_FUNCTION pair<KeyOutputIt, ValOutputIt> unique_by_key(
   ValOutputIt values_result,
   BinaryPred binary_pred)
 {
-  using size_type = typename iterator_traits<KeyInputIt>::difference_type;
+  using size_type = thrust::detail::it_difference_t<KeyInputIt>;
 
   size_type num_items = static_cast<size_type>(thrust::distance(keys_first, keys_last));
   pair<KeyOutputIt, ValOutputIt> result_end{};
@@ -273,7 +273,7 @@ pair<KeyOutputIt, ValOutputIt> _CCCL_HOST_DEVICE unique_by_key_copy(
   KeyOutputIt keys_result,
   ValOutputIt values_result)
 {
-  using key_type = typename iterator_traits<KeyInputIt>::value_type;
+  using key_type = thrust::detail::it_value_t<KeyInputIt>;
   return cuda_cub::unique_by_key_copy(
     policy, keys_first, keys_last, values_first, keys_result, values_result, equal_to<key_type>());
 }
@@ -298,7 +298,7 @@ template <class Derived, class KeyInputIt, class ValInputIt>
 pair<KeyInputIt, ValInputIt> _CCCL_HOST_DEVICE
 unique_by_key(execution_policy<Derived>& policy, KeyInputIt keys_first, KeyInputIt keys_last, ValInputIt values_first)
 {
-  using key_type = typename iterator_traits<KeyInputIt>::value_type;
+  using key_type = thrust::detail::it_value_t<KeyInputIt>;
   return cuda_cub::unique_by_key(policy, keys_first, keys_last, values_first, equal_to<key_type>());
 }
 
diff --git a/thrust/thrust/system/detail/generic/adjacent_difference.inl b/thrust/thrust/system/detail/generic/adjacent_difference.inl
index 275e44b16cd..23325102f71 100644
--- a/thrust/thrust/system/detail/generic/adjacent_difference.inl
+++ b/thrust/thrust/system/detail/generic/adjacent_difference.inl
@@ -44,7 +44,7 @@ template <typename DerivedPolicy, typename InputIterator, typename OutputIterato
 _CCCL_HOST_DEVICE OutputIterator adjacent_difference(
   thrust::execution_policy<DerivedPolicy>& exec, InputIterator first, InputIterator last, OutputIterator result)
 {
-  using InputType = typename thrust::iterator_traits<InputIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<InputIterator>;
   thrust::minus<InputType> binary_op;
 
   return thrust::adjacent_difference(exec, first, last, result, binary_op);
@@ -58,7 +58,7 @@ _CCCL_HOST_DEVICE OutputIterator adjacent_difference(
   OutputIterator result,
   BinaryFunction binary_op)
 {
-  using InputType = typename thrust::iterator_traits<InputIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<InputIterator>;
 
   if (first == last)
   {
diff --git a/thrust/thrust/system/detail/generic/binary_search.inl b/thrust/thrust/system/detail/generic/binary_search.inl
index ee049bbf1a5..bb201583aaf 100644
--- a/thrust/thrust/system/detail/generic/binary_search.inl
+++ b/thrust/thrust/system/detail/generic/binary_search.inl
@@ -58,7 +58,7 @@ namespace detail
 struct lbf
 {
   template <typename RandomAccessIterator, typename T, typename StrictWeakOrdering>
-  _CCCL_HOST_DEVICE typename thrust::iterator_traits<RandomAccessIterator>::difference_type
+  _CCCL_HOST_DEVICE thrust::detail::it_difference_t<RandomAccessIterator>
   operator()(RandomAccessIterator begin, RandomAccessIterator end, const T& value, StrictWeakOrdering comp)
   {
     return thrust::system::detail::generic::scalar::lower_bound(begin, end, value, comp) - begin;
@@ -68,7 +68,7 @@ struct lbf
 struct ubf
 {
   template <typename RandomAccessIterator, typename T, typename StrictWeakOrdering>
-  _CCCL_HOST_DEVICE typename thrust::iterator_traits<RandomAccessIterator>::difference_type
+  _CCCL_HOST_DEVICE thrust::detail::it_difference_t<RandomAccessIterator>
   operator()(RandomAccessIterator begin, RandomAccessIterator end, const T& value, StrictWeakOrdering comp)
   {
     return thrust::system::detail::generic::scalar::upper_bound(begin, end, value, comp) - begin;
@@ -224,7 +224,7 @@ _CCCL_HOST_DEVICE ForwardIterator lower_bound(
   const T& value,
   StrictWeakOrdering comp)
 {
-  using difference_type = typename thrust::iterator_traits<ForwardIterator>::difference_type;
+  using difference_type = thrust::detail::it_difference_t<ForwardIterator>;
 
   return begin + detail::binary_search<difference_type>(exec, begin, end, value, comp, detail::lbf());
 }
@@ -245,7 +245,7 @@ _CCCL_HOST_DEVICE ForwardIterator upper_bound(
   const T& value,
   StrictWeakOrdering comp)
 {
-  using difference_type = typename thrust::iterator_traits<ForwardIterator>::difference_type;
+  using difference_type = thrust::detail::it_difference_t<ForwardIterator>;
 
   return begin + detail::binary_search<difference_type>(exec, begin, end, value, comp, detail::ubf());
 }
diff --git a/thrust/thrust/system/detail/generic/copy_if.inl b/thrust/thrust/system/detail/generic/copy_if.inl
index 660913fedc2..e08f625fe8e 100644
--- a/thrust/thrust/system/detail/generic/copy_if.inl
+++ b/thrust/thrust/system/detail/generic/copy_if.inl
@@ -122,7 +122,7 @@ _CCCL_HOST_DEVICE OutputIterator copy_if(
   OutputIterator result,
   Predicate pred)
 {
-  using difference_type = typename thrust::iterator_traits<InputIterator1>::difference_type;
+  using difference_type = thrust::detail::it_difference_t<InputIterator1>;
 
   // empty sequence
   if (first == last)
diff --git a/thrust/thrust/system/detail/generic/count.h b/thrust/thrust/system/detail/generic/count.h
index dcef420932c..5c28590a54f 100644
--- a/thrust/thrust/system/detail/generic/count.h
+++ b/thrust/thrust/system/detail/generic/count.h
@@ -36,14 +36,14 @@ namespace generic
 {
 
 template <typename DerivedPolicy, typename InputIterator, typename EqualityComparable>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count(thrust::execution_policy<DerivedPolicy>& exec,
       InputIterator first,
       InputIterator last,
       const EqualityComparable& value);
 
 template <typename DerivedPolicy, typename InputIterator, typename Predicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count_if(thrust::execution_policy<DerivedPolicy>& exec, InputIterator first, InputIterator last, Predicate pred);
 
 } // end namespace generic
diff --git a/thrust/thrust/system/detail/generic/count.inl b/thrust/thrust/system/detail/generic/count.inl
index d5a1c2c7cd7..817478f83f3 100644
--- a/thrust/thrust/system/detail/generic/count.inl
+++ b/thrust/thrust/system/detail/generic/count.inl
@@ -61,7 +61,7 @@ struct count_if_transform
 }; // end count_if_transform
 
 template <typename DerivedPolicy, typename InputIterator, typename EqualityComparable>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count(thrust::execution_policy<DerivedPolicy>& exec,
       InputIterator first,
       InputIterator last,
@@ -73,11 +73,11 @@ count(thrust::execution_policy<DerivedPolicy>& exec,
 } // end count()
 
 template <typename DerivedPolicy, typename InputIterator, typename Predicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 count_if(thrust::execution_policy<DerivedPolicy>& exec, InputIterator first, InputIterator last, Predicate pred)
 {
-  using InputType = typename thrust::iterator_traits<InputIterator>::value_type;
-  using CountType = typename thrust::iterator_traits<InputIterator>::difference_type;
+  using InputType = thrust::detail::it_value_t<InputIterator>;
+  using CountType = thrust::detail::it_difference_t<InputIterator>;
 
   thrust::system::detail::generic::count_if_transform<InputType, Predicate, CountType> unary_op(pred);
   thrust::plus<CountType> binary_op;
diff --git a/thrust/thrust/system/detail/generic/distance.h b/thrust/thrust/system/detail/generic/distance.h
index 30c79a2387b..f68b14cf7b8 100644
--- a/thrust/thrust/system/detail/generic/distance.h
+++ b/thrust/thrust/system/detail/generic/distance.h
@@ -36,7 +36,7 @@ namespace generic
 {
 
 template <typename InputIterator>
-inline _CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+inline _CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 distance(InputIterator first, InputIterator last);
 
 } // end namespace generic
diff --git a/thrust/thrust/system/detail/generic/distance.inl b/thrust/thrust/system/detail/generic/distance.inl
index e12a0ea6879..52b8eb282b0 100644
--- a/thrust/thrust/system/detail/generic/distance.inl
+++ b/thrust/thrust/system/detail/generic/distance.inl
@@ -40,10 +40,10 @@ namespace detail
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename InputIterator>
-inline _CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+inline _CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 distance(InputIterator first, InputIterator last, thrust::incrementable_traversal_tag)
 {
-  typename thrust::iterator_traits<InputIterator>::difference_type result(0);
+  thrust::detail::it_difference_t<InputIterator> result(0);
 
   while (first != last)
   {
@@ -56,7 +56,7 @@ distance(InputIterator first, InputIterator last, thrust::incrementable_traversa
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename InputIterator>
-inline _CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
+inline _CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator>
 distance(InputIterator first, InputIterator last, thrust::random_access_traversal_tag)
 {
   return last - first;
@@ -66,8 +66,7 @@ distance(InputIterator first, InputIterator last, thrust::random_access_traversa
 
 _CCCL_EXEC_CHECK_DISABLE
 template <typename InputIterator>
-inline _CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::difference_type
-distance(InputIterator first, InputIterator last)
+inline _CCCL_HOST_DEVICE thrust::detail::it_difference_t<InputIterator> distance(InputIterator first, InputIterator last)
 {
   // dispatch on iterator traversal
   return thrust::system::detail::generic::detail::distance(
diff --git a/thrust/thrust/system/detail/generic/extrema.inl b/thrust/thrust/system/detail/generic/extrema.inl
index 0683250e92b..a64fd830dc2 100644
--- a/thrust/thrust/system/detail/generic/extrema.inl
+++ b/thrust/thrust/system/detail/generic/extrema.inl
@@ -159,7 +159,7 @@ template <typename DerivedPolicy, typename ForwardIterator>
 _CCCL_HOST_DEVICE ForwardIterator
 min_element(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last)
 {
-  using value_type = typename thrust::iterator_value<ForwardIterator>::type;
+  using value_type = thrust::detail::it_value_t<ForwardIterator>;
 
   return thrust::min_element(exec, first, last, thrust::less<value_type>());
 } // end min_element()
@@ -173,8 +173,8 @@ _CCCL_HOST_DEVICE ForwardIterator min_element(
     return last;
   }
 
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
-  using IndexType = typename thrust::iterator_traits<ForwardIterator>::difference_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
+  using IndexType = thrust::detail::it_difference_t<ForwardIterator>;
 
   thrust::tuple<InputType, IndexType> result = thrust::reduce(
     exec,
@@ -190,7 +190,7 @@ template <typename DerivedPolicy, typename ForwardIterator>
 _CCCL_HOST_DEVICE ForwardIterator
 max_element(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last)
 {
-  using value_type = typename thrust::iterator_value<ForwardIterator>::type;
+  using value_type = thrust::detail::it_value_t<ForwardIterator>;
 
   return thrust::max_element(exec, first, last, thrust::less<value_type>());
 } // end max_element()
@@ -204,8 +204,8 @@ _CCCL_HOST_DEVICE ForwardIterator max_element(
     return last;
   }
 
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
-  using IndexType = typename thrust::iterator_traits<ForwardIterator>::difference_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
+  using IndexType = thrust::detail::it_difference_t<ForwardIterator>;
 
   thrust::tuple<InputType, IndexType> result = thrust::reduce(
     exec,
@@ -221,7 +221,7 @@ template <typename DerivedPolicy, typename ForwardIterator>
 _CCCL_HOST_DEVICE thrust::pair<ForwardIterator, ForwardIterator>
 minmax_element(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last)
 {
-  using value_type = typename thrust::iterator_value<ForwardIterator>::type;
+  using value_type = thrust::detail::it_value_t<ForwardIterator>;
 
   return thrust::minmax_element(exec, first, last, thrust::less<value_type>());
 } // end minmax_element()
@@ -235,8 +235,8 @@ _CCCL_HOST_DEVICE thrust::pair<ForwardIterator, ForwardIterator> minmax_element(
     return thrust::make_pair(last, last);
   }
 
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
-  using IndexType = typename thrust::iterator_traits<ForwardIterator>::difference_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
+  using IndexType = thrust::detail::it_difference_t<ForwardIterator>;
 
   thrust::tuple<thrust::tuple<InputType, IndexType>, thrust::tuple<InputType, IndexType>> result =
     thrust::transform_reduce(
diff --git a/thrust/thrust/system/detail/generic/find.inl b/thrust/thrust/system/detail/generic/find.inl
index 09ec0b8c59e..f77bb3883c1 100644
--- a/thrust/thrust/system/detail/generic/find.inl
+++ b/thrust/thrust/system/detail/generic/find.inl
@@ -79,7 +79,7 @@ template <typename DerivedPolicy, typename InputIterator, typename Predicate>
 _CCCL_HOST_DEVICE InputIterator
 find_if(thrust::execution_policy<DerivedPolicy>& exec, InputIterator first, InputIterator last, Predicate pred)
 {
-  using difference_type = typename thrust::iterator_traits<InputIterator>::difference_type;
+  using difference_type = thrust::detail::it_difference_t<InputIterator>;
   using result_type     = typename thrust::tuple<bool, difference_type>;
 
   // empty sequence
diff --git a/thrust/thrust/system/detail/generic/generate.inl b/thrust/thrust/system/detail/generic/generate.inl
index cc72f7ca583..7705980b2e3 100644
--- a/thrust/thrust/system/detail/generic/generate.inl
+++ b/thrust/thrust/system/detail/generic/generate.inl
@@ -55,8 +55,7 @@ generate(thrust::execution_policy<ExecutionPolicy>& exec, ForwardIterator first,
   // nice solution that validates the const_cast and doesn't take away any
   // functionality.
   THRUST_STATIC_ASSERT_MSG(
-    !::cuda::std::is_const<
-      ::cuda::std::remove_reference_t<typename thrust::iterator_traits<ForwardIterator>::reference>>::value,
+    !::cuda::std::is_const<::cuda::std::remove_reference_t<thrust::detail::it_reference_t<ForwardIterator>>>::value,
     "generating to `const` iterators is not allowed");
   thrust::for_each(exec, first, last, typename thrust::detail::generate_functor<ExecutionPolicy, Generator>::type(gen));
 } // end generate()
@@ -77,8 +76,7 @@ generate_n(thrust::execution_policy<ExecutionPolicy>& exec, OutputIterator first
   // nice solution that validates the const_cast and doesn't take away any
   // functionality.
   THRUST_STATIC_ASSERT_MSG(
-    !::cuda::std::is_const<
-      ::cuda::std::remove_reference_t<typename thrust::iterator_traits<OutputIterator>::reference>>::value,
+    !::cuda::std::is_const<::cuda::std::remove_reference_t<thrust::detail::it_reference_t<OutputIterator>>>::value,
     "generating to `const` iterators is not allowed");
   return thrust::for_each_n(
     exec, first, n, typename thrust::detail::generate_functor<ExecutionPolicy, Generator>::type(gen));
diff --git a/thrust/thrust/system/detail/generic/merge.inl b/thrust/thrust/system/detail/generic/merge.inl
index 67336eb87ca..17700a605f1 100644
--- a/thrust/thrust/system/detail/generic/merge.inl
+++ b/thrust/thrust/system/detail/generic/merge.inl
@@ -68,7 +68,7 @@ _CCCL_HOST_DEVICE OutputIterator merge(
   InputIterator2 last2,
   OutputIterator result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::merge(exec, first1, last1, first2, last2, result, thrust::less<value_type>());
 } // end merge()
 
@@ -135,7 +135,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> merge_by_key(
   OutputIterator1 keys_result,
   OutputIterator2 values_result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::merge_by_key(
     exec,
     keys_first1,
diff --git a/thrust/thrust/system/detail/generic/partition.inl b/thrust/thrust/system/detail/generic/partition.inl
index 67a1f9c2a44..90826c63554 100644
--- a/thrust/thrust/system/detail/generic/partition.inl
+++ b/thrust/thrust/system/detail/generic/partition.inl
@@ -49,13 +49,13 @@ template <typename DerivedPolicy, typename ForwardIterator, typename Predicate>
 _CCCL_HOST_DEVICE ForwardIterator stable_partition(
   thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last, Predicate pred)
 {
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
 
   // copy input to temp buffer
   thrust::detail::temporary_array<InputType, DerivedPolicy> temp(exec, first, last);
 
   // count the size of the true partition
-  typename thrust::iterator_difference<ForwardIterator>::type num_true = thrust::count_if(exec, first, last, pred);
+  thrust::detail::it_difference_t<ForwardIterator> num_true = thrust::count_if(exec, first, last, pred);
 
   // point to the beginning of the false partition
   ForwardIterator out_false = first;
@@ -72,7 +72,7 @@ _CCCL_HOST_DEVICE ForwardIterator stable_partition(
   InputIterator stencil,
   Predicate pred)
 {
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
 
   // copy input to temp buffer
   thrust::detail::temporary_array<InputType, DerivedPolicy> temp(exec, first, last);
@@ -80,8 +80,7 @@ _CCCL_HOST_DEVICE ForwardIterator stable_partition(
   // count the size of the true partition
   InputIterator stencil_last = stencil;
   thrust::advance(stencil_last, temp.size());
-  typename thrust::iterator_difference<InputIterator>::type num_true =
-    thrust::count_if(exec, stencil, stencil_last, pred);
+  thrust::detail::it_difference_t<InputIterator> num_true = thrust::count_if(exec, stencil, stencil_last, pred);
 
   // point to the beginning of the false partition
   ForwardIterator out_false = first;
diff --git a/thrust/thrust/system/detail/generic/reduce.h b/thrust/thrust/system/detail/generic/reduce.h
index b1fe9a51ed0..0fd12dca12a 100644
--- a/thrust/thrust/system/detail/generic/reduce.h
+++ b/thrust/thrust/system/detail/generic/reduce.h
@@ -37,7 +37,7 @@ namespace generic
 {
 
 template <typename DerivedPolicy, typename InputIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::value_type
+_CCCL_HOST_DEVICE thrust::detail::it_value_t<InputIterator>
 reduce(thrust::execution_policy<DerivedPolicy>& exec, InputIterator first, InputIterator last);
 
 template <typename DerivedPolicy, typename InputIterator, typename T>
diff --git a/thrust/thrust/system/detail/generic/reduce.inl b/thrust/thrust/system/detail/generic/reduce.inl
index 7164103141f..6401065e192 100644
--- a/thrust/thrust/system/detail/generic/reduce.inl
+++ b/thrust/thrust/system/detail/generic/reduce.inl
@@ -41,10 +41,10 @@ namespace generic
 {
 
 template <typename ExecutionPolicy, typename InputIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<InputIterator>::value_type
+_CCCL_HOST_DEVICE thrust::detail::it_value_t<InputIterator>
 reduce(thrust::execution_policy<ExecutionPolicy>& exec, InputIterator first, InputIterator last)
 {
-  using InputType = typename thrust::iterator_value<InputIterator>::type;
+  using InputType = thrust::detail::it_value_t<InputIterator>;
 
   // use InputType(0) as init by default
   return thrust::reduce(exec, first, last, InputType(0));
diff --git a/thrust/thrust/system/detail/generic/reduce_by_key.inl b/thrust/thrust/system/detail/generic/reduce_by_key.inl
index 2d0d73859f0..a24735b5ea5 100644
--- a/thrust/thrust/system/detail/generic/reduce_by_key.inl
+++ b/thrust/thrust/system/detail/generic/reduce_by_key.inl
@@ -37,6 +37,8 @@
 #include <thrust/scatter.h>
 #include <thrust/transform.h>
 
+#include <cuda/std/iterator>
+
 #include <limits>
 
 THRUST_NAMESPACE_BEGIN
@@ -86,12 +88,12 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> reduce_by_key(
   BinaryPredicate binary_pred,
   BinaryFunction binary_op)
 {
-  using difference_type = typename thrust::iterator_traits<InputIterator1>::difference_type;
+  using difference_type = thrust::detail::it_difference_t<InputIterator1>;
 
   using FlagType = unsigned int; // TODO use difference_type
 
   // Use the input iterator's value type per https://wg21.link/P0571
-  using ValueType = typename thrust::iterator_value<InputIterator2>::type;
+  using ValueType = thrust::detail::it_value_t<InputIterator2>;
 
   if (keys_first == keys_last)
   {
@@ -152,7 +154,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> reduce_by_key(
   OutputIterator1 keys_output,
   OutputIterator2 values_output)
 {
-  using KeyType = typename thrust::iterator_value<InputIterator1>::type;
+  using KeyType = thrust::detail::it_value_t<InputIterator1>;
 
   // use equal_to<KeyType> as default BinaryPredicate
   return thrust::reduce_by_key(
@@ -174,9 +176,11 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> reduce_by_key(
   OutputIterator2 values_output,
   BinaryPredicate binary_pred)
 {
-  using T = typename thrust::detail::eval_if<thrust::detail::is_output_iterator<OutputIterator2>::value,
-                                             thrust::iterator_value<InputIterator2>,
-                                             thrust::iterator_value<OutputIterator2>>::type;
+  using T = ::cuda::std::
+
+    _If<thrust::detail::is_output_iterator<OutputIterator2>,
+        thrust::detail::it_value_t<InputIterator2>,
+        thrust::detail::it_value_t<OutputIterator2>>;
 
   // use plus<T> as default BinaryFunction
   return thrust::reduce_by_key(
diff --git a/thrust/thrust/system/detail/generic/remove.inl b/thrust/thrust/system/detail/generic/remove.inl
index cf33d436987..89373df1e12 100644
--- a/thrust/thrust/system/detail/generic/remove.inl
+++ b/thrust/thrust/system/detail/generic/remove.inl
@@ -68,7 +68,7 @@ template <typename DerivedPolicy, typename ForwardIterator, typename Predicate>
 _CCCL_HOST_DEVICE ForwardIterator
 remove_if(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last, Predicate pred)
 {
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
 
   // create temporary storage for an intermediate result
   thrust::detail::temporary_array<InputType, DerivedPolicy> temp(exec, first, last);
@@ -85,7 +85,7 @@ _CCCL_HOST_DEVICE ForwardIterator remove_if(
   InputIterator stencil,
   Predicate pred)
 {
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
 
   // create temporary storage for an intermediate result
   thrust::detail::temporary_array<InputType, DerivedPolicy> temp(exec, first, last);
diff --git a/thrust/thrust/system/detail/generic/replace.inl b/thrust/thrust/system/detail/generic/replace.inl
index 5c09b58c311..c7dfb9019c6 100644
--- a/thrust/thrust/system/detail/generic/replace.inl
+++ b/thrust/thrust/system/detail/generic/replace.inl
@@ -96,7 +96,7 @@ _CCCL_HOST_DEVICE OutputIterator replace_copy_if(
   Predicate pred,
   const T& new_value)
 {
-  using OutputType = typename thrust::iterator_traits<OutputIterator>::value_type;
+  using OutputType = thrust::detail::it_value_t<OutputIterator>;
 
   detail::new_value_if<Predicate, T, OutputType> op(pred, new_value);
   return thrust::transform(exec, first, last, result, op);
@@ -117,7 +117,7 @@ _CCCL_HOST_DEVICE OutputIterator replace_copy_if(
   Predicate pred,
   const T& new_value)
 {
-  using OutputType = typename thrust::iterator_traits<OutputIterator>::value_type;
+  using OutputType = thrust::detail::it_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/reverse.inl b/thrust/thrust/system/detail/generic/reverse.inl
index acc184152c3..19d119029e8 100644
--- a/thrust/thrust/system/detail/generic/reverse.inl
+++ b/thrust/thrust/system/detail/generic/reverse.inl
@@ -45,7 +45,7 @@ template <typename ExecutionPolicy, typename BidirectionalIterator>
 _CCCL_HOST_DEVICE void
 reverse(thrust::execution_policy<ExecutionPolicy>& exec, BidirectionalIterator first, BidirectionalIterator last)
 {
-  using difference_type = typename thrust::iterator_difference<BidirectionalIterator>::type;
+  using difference_type = thrust::detail::it_difference_t<BidirectionalIterator>;
 
   // find the midpoint of [first,last)
   difference_type N = thrust::distance(first, last);
diff --git a/thrust/thrust/system/detail/generic/scalar/binary_search.inl b/thrust/thrust/system/detail/generic/scalar/binary_search.inl
index 52e1a637f23..066076bd4bb 100644
--- a/thrust/thrust/system/detail/generic/scalar/binary_search.inl
+++ b/thrust/thrust/system/detail/generic/scalar/binary_search.inl
@@ -73,7 +73,7 @@ template <typename RandomAccessIterator, typename T, typename BinaryPredicate>
 _CCCL_HOST_DEVICE RandomAccessIterator
 lower_bound(RandomAccessIterator first, RandomAccessIterator last, const T& val, BinaryPredicate comp)
 {
-  typename thrust::iterator_difference<RandomAccessIterator>::type n = last - first;
+  thrust::detail::it_difference_t<RandomAccessIterator> n = last - first;
   return lower_bound_n(first, n, val, comp);
 }
 
@@ -105,7 +105,7 @@ template <typename RandomAccessIterator, typename T, typename BinaryPredicate>
 _CCCL_HOST_DEVICE RandomAccessIterator
 upper_bound(RandomAccessIterator first, RandomAccessIterator last, const T& val, BinaryPredicate comp)
 {
-  typename thrust::iterator_difference<RandomAccessIterator>::type n = last - first;
+  thrust::detail::it_difference_t<RandomAccessIterator> n = last - first;
   return upper_bound_n(first, n, val, comp);
 }
 
diff --git a/thrust/thrust/system/detail/generic/scan.inl b/thrust/thrust/system/detail/generic/scan.inl
index 05dca4bb5a6..dc709d7700b 100644
--- a/thrust/thrust/system/detail/generic/scan.inl
+++ b/thrust/thrust/system/detail/generic/scan.inl
@@ -27,7 +27,6 @@
 #endif // no system header
 #include <thrust/detail/static_assert.h>
 #include <thrust/detail/type_traits.h>
-#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/scan.h>
@@ -54,7 +53,7 @@ _CCCL_HOST_DEVICE OutputIterator exclusive_scan(
   thrust::execution_policy<ExecutionPolicy>& exec, InputIterator first, InputIterator last, OutputIterator result)
 {
   // Use the input iterator's value type per https://wg21.link/P0571
-  using ValueType = typename thrust::iterator_value<InputIterator>::type;
+  using ValueType = thrust::detail::it_value_t<InputIterator>;
   return thrust::exclusive_scan(exec, first, last, result, ValueType{});
 } // end exclusive_scan()
 
diff --git a/thrust/thrust/system/detail/generic/scan_by_key.inl b/thrust/thrust/system/detail/generic/scan_by_key.inl
index fe4a68e13e0..f545ba73545 100644
--- a/thrust/thrust/system/detail/generic/scan_by_key.inl
+++ b/thrust/thrust/system/detail/generic/scan_by_key.inl
@@ -109,7 +109,7 @@ _CCCL_HOST_DEVICE OutputIterator inclusive_scan_by_key(
   BinaryPredicate binary_pred,
   AssociativeOperator binary_op)
 {
-  using OutputType   = typename thrust::iterator_traits<InputIterator2>::value_type;
+  using OutputType   = thrust::detail::it_value_t<InputIterator2>;
   using HeadFlagType = std::uint8_t;
 
   const size_t n = last1 - first1;
@@ -145,7 +145,7 @@ _CCCL_HOST_DEVICE OutputIterator exclusive_scan_by_key(
   InputIterator2 first2,
   OutputIterator result)
 {
-  using InitType = typename thrust::iterator_traits<InputIterator2>::value_type;
+  using InitType = thrust::detail::it_value_t<InputIterator2>;
   return thrust::exclusive_scan_by_key(exec, first1, last1, first2, result, InitType{});
 }
 
diff --git a/thrust/thrust/system/detail/generic/sequence.inl b/thrust/thrust/system/detail/generic/sequence.inl
index 2c26725ebbd..2cd044096cd 100644
--- a/thrust/thrust/system/detail/generic/sequence.inl
+++ b/thrust/thrust/system/detail/generic/sequence.inl
@@ -41,7 +41,7 @@ template <typename DerivedPolicy, typename ForwardIterator>
 _CCCL_HOST_DEVICE void
 sequence(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last)
 {
-  using T = typename thrust::iterator_traits<ForwardIterator>::value_type;
+  using T = thrust::detail::it_value_t<ForwardIterator>;
 
   thrust::sequence(exec, first, last, T(0));
 } // end sequence()
diff --git a/thrust/thrust/system/detail/generic/set_operations.inl b/thrust/thrust/system/detail/generic/set_operations.inl
index 45efd79001e..9f321288349 100644
--- a/thrust/thrust/system/detail/generic/set_operations.inl
+++ b/thrust/thrust/system/detail/generic/set_operations.inl
@@ -50,7 +50,7 @@ _CCCL_HOST_DEVICE OutputIterator set_difference(
   InputIterator2 last2,
   OutputIterator result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::set_difference(exec, first1, last1, first2, last2, result, thrust::less<value_type>());
 } // end set_difference()
 
@@ -72,7 +72,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> set_difference_
   OutputIterator1 keys_result,
   OutputIterator2 values_result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::set_difference_by_key(
     exec,
     keys_first1,
@@ -140,7 +140,7 @@ _CCCL_HOST_DEVICE OutputIterator set_intersection(
   InputIterator2 last2,
   OutputIterator result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::set_intersection(exec, first1, last1, first2, last2, result, thrust::less<value_type>());
 } // end set_intersection()
 
@@ -160,7 +160,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> set_intersectio
   OutputIterator1 keys_result,
   OutputIterator2 values_result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::set_intersection_by_key(
     exec,
     keys_first1,
@@ -191,7 +191,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> set_intersectio
   OutputIterator2 values_result,
   StrictWeakOrdering comp)
 {
-  using value_type1       = typename thrust::iterator_value<InputIterator3>::type;
+  using value_type1       = thrust::detail::it_value_t<InputIterator3>;
   using constant_iterator = thrust::constant_iterator<value_type1>;
 
   using iterator_tuple1 = thrust::tuple<InputIterator1, InputIterator3>;
@@ -232,7 +232,7 @@ _CCCL_HOST_DEVICE OutputIterator set_symmetric_difference(
   InputIterator2 last2,
   OutputIterator result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::set_symmetric_difference(exec, first1, last1, first2, last2, result, thrust::less<value_type>());
 } // end set_symmetric_difference()
 
@@ -254,7 +254,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> set_symmetric_d
   OutputIterator1 keys_result,
   OutputIterator2 values_result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::set_symmetric_difference_by_key(
     exec,
     keys_first1,
@@ -323,7 +323,7 @@ _CCCL_HOST_DEVICE OutputIterator set_union(
   InputIterator2 last2,
   OutputIterator result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::set_union(exec, first1, last1, first2, last2, result, thrust::less<value_type>());
 } // end set_union()
 
@@ -345,7 +345,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> set_union_by_ke
   OutputIterator1 keys_result,
   OutputIterator2 values_result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator1>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator1>;
   return thrust::set_union_by_key(
     exec,
     keys_first1,
diff --git a/thrust/thrust/system/detail/generic/shuffle.inl b/thrust/thrust/system/detail/generic/shuffle.inl
index 16a097407b0..70d96de830d 100644
--- a/thrust/thrust/system/detail/generic/shuffle.inl
+++ b/thrust/thrust/system/detail/generic/shuffle.inl
@@ -171,7 +171,7 @@ template <typename ExecutionPolicy, typename RandomIterator, typename URBG>
 _CCCL_HOST_DEVICE void
 shuffle(thrust::execution_policy<ExecutionPolicy>& exec, RandomIterator first, RandomIterator last, URBG&& g)
 {
-  using InputType = typename thrust::iterator_value_t<RandomIterator>;
+  using InputType = typename thrust::detail::it_value_t<RandomIterator>;
 
   // copy input to temp buffer
   thrust::detail::temporary_array<InputType, ExecutionPolicy> temp(exec, first, last);
diff --git a/thrust/thrust/system/detail/generic/sort.inl b/thrust/thrust/system/detail/generic/sort.inl
index 1ea438868a9..cb1d0474e27 100644
--- a/thrust/thrust/system/detail/generic/sort.inl
+++ b/thrust/thrust/system/detail/generic/sort.inl
@@ -46,7 +46,7 @@ template <typename DerivedPolicy, typename RandomAccessIterator>
 _CCCL_HOST_DEVICE void
 sort(thrust::execution_policy<DerivedPolicy>& exec, RandomAccessIterator first, RandomAccessIterator last)
 {
-  using value_type = typename thrust::iterator_value<RandomAccessIterator>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator>;
   thrust::sort(exec, first, last, thrust::less<value_type>());
 } // end sort()
 
@@ -68,7 +68,7 @@ _CCCL_HOST_DEVICE void sort_by_key(
   RandomAccessIterator1 keys_last,
   RandomAccessIterator2 values_first)
 {
-  using value_type = typename thrust::iterator_value<RandomAccessIterator1>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator1>;
   thrust::sort_by_key(exec, keys_first, keys_last, values_first, thrust::less<value_type>());
 } // end sort_by_key()
 
@@ -91,7 +91,7 @@ template <typename DerivedPolicy, typename RandomAccessIterator>
 _CCCL_HOST_DEVICE void
 stable_sort(thrust::execution_policy<DerivedPolicy>& exec, RandomAccessIterator first, RandomAccessIterator last)
 {
-  using value_type = typename thrust::iterator_value<RandomAccessIterator>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator>;
   thrust::stable_sort(exec, first, last, thrust::less<value_type>());
 } // end stable_sort()
 
@@ -102,7 +102,7 @@ _CCCL_HOST_DEVICE void stable_sort_by_key(
   RandomAccessIterator1 keys_last,
   RandomAccessIterator2 values_first)
 {
-  using value_type = typename iterator_value<RandomAccessIterator1>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator1>;
   thrust::stable_sort_by_key(exec, keys_first, keys_last, values_first, thrust::less<value_type>());
 } // end stable_sort_by_key()
 
@@ -124,7 +124,7 @@ template <typename DerivedPolicy, typename ForwardIterator>
 _CCCL_HOST_DEVICE ForwardIterator
 is_sorted_until(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last)
 {
-  using InputType = typename thrust::iterator_value<ForwardIterator>::type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
 
   return thrust::is_sorted_until(exec, first, last, thrust::less<InputType>());
 } // end is_sorted_until()
diff --git a/thrust/thrust/system/detail/generic/tabulate.inl b/thrust/thrust/system/detail/generic/tabulate.inl
index 063e8b6e85c..6ce1c9aa513 100644
--- a/thrust/thrust/system/detail/generic/tabulate.inl
+++ b/thrust/thrust/system/detail/generic/tabulate.inl
@@ -43,7 +43,7 @@ template <typename DerivedPolicy, typename ForwardIterator, typename UnaryOperat
 _CCCL_HOST_DEVICE void tabulate(
   thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last, UnaryOperation unary_op)
 {
-  using difference_type = typename iterator_difference<ForwardIterator>::type;
+  using difference_type = thrust::detail::it_difference_t<ForwardIterator>;
 
   // by default, counting_iterator uses a 64b difference_type on 32b platforms to avoid overflowing its counter.
   // this causes problems when a zip_iterator is created in transform's implementation -- ForwardIterator is
diff --git a/thrust/thrust/system/detail/generic/transform_scan.inl b/thrust/thrust/system/detail/generic/transform_scan.inl
index fbd5f24b21e..f244d5bfab0 100644
--- a/thrust/thrust/system/detail/generic/transform_scan.inl
+++ b/thrust/thrust/system/detail/generic/transform_scan.inl
@@ -26,7 +26,6 @@
 #  pragma system_header
 #endif // no system header
 #include <thrust/detail/type_traits.h>
-#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/system/detail/generic/transform_scan.h>
@@ -55,7 +54,7 @@ _CCCL_HOST_DEVICE OutputIterator transform_inclusive_scan(
   BinaryFunction binary_op)
 {
   // Use the input iterator's value type per https://wg21.link/P0571
-  using InputType  = typename thrust::iterator_value<InputIterator>::type;
+  using InputType  = thrust::detail::it_value_t<InputIterator>;
   using ResultType = thrust::detail::invoke_result_t<UnaryFunction, InputType>;
   using ValueType  = ::cuda::std::remove_cvref_t<ResultType>;
 
@@ -80,7 +79,7 @@ _CCCL_HOST_DEVICE OutputIterator transform_inclusive_scan(
   InitialValueType init,
   BinaryFunction binary_op)
 {
-  using InputType  = typename thrust::iterator_value<InputIterator>::type;
+  using InputType  = thrust::detail::it_value_t<InputIterator>;
   using ResultType = thrust::detail::invoke_result_t<UnaryFunction, InputType>;
   using ValueType  = ::cuda::std::remove_cvref_t<ResultType>;
 
diff --git a/thrust/thrust/system/detail/generic/uninitialized_copy.inl b/thrust/thrust/system/detail/generic/uninitialized_copy.inl
index 931d023e2dc..377f529168c 100644
--- a/thrust/thrust/system/detail/generic/uninitialized_copy.inl
+++ b/thrust/thrust/system/detail/generic/uninitialized_copy.inl
@@ -72,12 +72,12 @@ _CCCL_HOST_DEVICE ForwardIterator uninitialized_copy(
   ZipIterator end   = begin;
 
   // get a zip_iterator pointing to the end
-  const typename thrust::iterator_difference<InputIterator>::type n = thrust::distance(first, last);
+  const thrust::detail::it_difference_t<InputIterator> n = thrust::distance(first, last);
   thrust::advance(end, n);
 
   // create a functor
-  using InputType  = typename iterator_traits<InputIterator>::value_type;
-  using OutputType = typename iterator_traits<ForwardIterator>::value_type;
+  using InputType  = thrust::detail::it_value_t<InputIterator>;
+  using OutputType = thrust::detail::it_value_t<ForwardIterator>;
 
   detail::uninitialized_copy_functor<InputType, OutputType> f;
 
@@ -116,8 +116,8 @@ _CCCL_HOST_DEVICE ForwardIterator uninitialized_copy_n(
   ZipIterator zipped_first = thrust::make_zip_iterator(thrust::make_tuple(first, result));
 
   // create a functor
-  using InputType  = typename iterator_traits<InputIterator>::value_type;
-  using OutputType = typename iterator_traits<ForwardIterator>::value_type;
+  using InputType  = thrust::detail::it_value_t<InputIterator>;
+  using OutputType = thrust::detail::it_value_t<ForwardIterator>;
 
   detail::uninitialized_copy_functor<InputType, OutputType> f;
 
@@ -146,7 +146,7 @@ template <typename ExecutionPolicy, typename InputIterator, typename ForwardIter
 _CCCL_HOST_DEVICE ForwardIterator uninitialized_copy(
   thrust::execution_policy<ExecutionPolicy>& exec, InputIterator first, InputIterator last, ForwardIterator result)
 {
-  using ResultType = typename iterator_traits<ForwardIterator>::value_type;
+  using ResultType = thrust::detail::it_value_t<ForwardIterator>;
 
   using ResultTypeHasTrivialCopyConstructor = typename ::cuda::std::is_trivially_copy_constructible<ResultType>::type;
 
@@ -158,7 +158,7 @@ template <typename ExecutionPolicy, typename InputIterator, typename Size, typen
 _CCCL_HOST_DEVICE ForwardIterator uninitialized_copy_n(
   thrust::execution_policy<ExecutionPolicy>& exec, InputIterator first, Size n, ForwardIterator result)
 {
-  using ResultType = typename iterator_traits<ForwardIterator>::value_type;
+  using ResultType = thrust::detail::it_value_t<ForwardIterator>;
 
   using ResultTypeHasTrivialCopyConstructor = typename ::cuda::std::is_trivially_copy_constructible<ResultType>::type;
 
diff --git a/thrust/thrust/system/detail/generic/uninitialized_fill.inl b/thrust/thrust/system/detail/generic/uninitialized_fill.inl
index e8717c23966..02871da5dd6 100644
--- a/thrust/thrust/system/detail/generic/uninitialized_fill.inl
+++ b/thrust/thrust/system/detail/generic/uninitialized_fill.inl
@@ -60,7 +60,7 @@ _CCCL_HOST_DEVICE void uninitialized_fill(
   const T& x,
   thrust::detail::false_type) // ::cuda::std::is_trivially_copy_constructible
 {
-  using ValueType = typename iterator_traits<ForwardIterator>::value_type;
+  using ValueType = thrust::detail::it_value_t<ForwardIterator>;
 
   thrust::for_each(exec, first, last, thrust::detail::uninitialized_fill_functor<ValueType>(x));
 } // end uninitialized_fill()
@@ -84,7 +84,7 @@ _CCCL_HOST_DEVICE ForwardIterator uninitialized_fill_n(
   const T& x,
   thrust::detail::false_type) // ::cuda::std::is_trivially_copy_constructible
 {
-  using ValueType = typename iterator_traits<ForwardIterator>::value_type;
+  using ValueType = thrust::detail::it_value_t<ForwardIterator>;
 
   return thrust::for_each_n(exec, first, n, thrust::detail::uninitialized_fill_functor<ValueType>(x));
 } // end uninitialized_fill()
@@ -95,7 +95,7 @@ template <typename DerivedPolicy, typename ForwardIterator, typename T>
 _CCCL_HOST_DEVICE void uninitialized_fill(
   thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last, const T& x)
 {
-  using ValueType = typename iterator_traits<ForwardIterator>::value_type;
+  using ValueType = thrust::detail::it_value_t<ForwardIterator>;
 
   using ValueTypeHasTrivialCopyConstructor = ::cuda::std::is_trivially_copy_constructible<ValueType>;
 
@@ -107,7 +107,7 @@ template <typename DerivedPolicy, typename ForwardIterator, typename Size, typen
 _CCCL_HOST_DEVICE ForwardIterator
 uninitialized_fill_n(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, Size n, const T& x)
 {
-  using ValueType = typename iterator_traits<ForwardIterator>::value_type;
+  using ValueType = thrust::detail::it_value_t<ForwardIterator>;
 
   using ValueTypeHasTrivialCopyConstructor = ::cuda::std::is_trivially_copy_constructible<ValueType>;
 
diff --git a/thrust/thrust/system/detail/generic/unique.h b/thrust/thrust/system/detail/generic/unique.h
index fc0f01311cb..e39470cdd74 100644
--- a/thrust/thrust/system/detail/generic/unique.h
+++ b/thrust/thrust/system/detail/generic/unique.h
@@ -59,11 +59,11 @@ _CCCL_HOST_DEVICE OutputIterator unique_copy(
   BinaryPredicate binary_pred);
 
 template <typename DerivedPolicy, typename ForwardIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator>
 unique_count(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last);
 
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator> unique_count(
   thrust::execution_policy<DerivedPolicy>& exec,
   ForwardIterator first,
   ForwardIterator last,
diff --git a/thrust/thrust/system/detail/generic/unique.inl b/thrust/thrust/system/detail/generic/unique.inl
index c43581f95f0..aac338d1d0c 100644
--- a/thrust/thrust/system/detail/generic/unique.inl
+++ b/thrust/thrust/system/detail/generic/unique.inl
@@ -49,7 +49,7 @@ template <typename DerivedPolicy, typename ForwardIterator>
 _CCCL_HOST_DEVICE ForwardIterator
 unique(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last)
 {
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
 
   return thrust::unique(exec, first, last, thrust::equal_to<InputType>());
 } // end unique()
@@ -61,7 +61,7 @@ _CCCL_HOST_DEVICE ForwardIterator unique(
   ForwardIterator last,
   BinaryPredicate binary_pred)
 {
-  using InputType = typename thrust::iterator_traits<ForwardIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<ForwardIterator>;
 
   thrust::detail::temporary_array<InputType, DerivedPolicy> input(exec, first, last);
 
@@ -72,7 +72,7 @@ template <typename DerivedPolicy, typename InputIterator, typename OutputIterato
 _CCCL_HOST_DEVICE OutputIterator unique_copy(
   thrust::execution_policy<DerivedPolicy>& exec, InputIterator first, InputIterator last, OutputIterator output)
 {
-  using value_type = typename thrust::iterator_value<InputIterator>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator>;
   return thrust::unique_copy(exec, first, last, output, thrust::equal_to<value_type>());
 } // end unique_copy()
 
@@ -92,7 +92,7 @@ _CCCL_HOST_DEVICE OutputIterator unique_copy(
 } // end unique_copy()
 
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator> unique_count(
   thrust::execution_policy<DerivedPolicy>& exec,
   ForwardIterator first,
   ForwardIterator last,
@@ -106,10 +106,10 @@ _CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_
 } // end unique_copy()
 
 template <typename DerivedPolicy, typename ForwardIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator>
 unique_count(thrust::execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last)
 {
-  using value_type = typename thrust::iterator_value<ForwardIterator>::type;
+  using value_type = thrust::detail::it_value_t<ForwardIterator>;
   return thrust::unique_count(exec, first, last, thrust::equal_to<value_type>());
 } // end unique_copy()
 
diff --git a/thrust/thrust/system/detail/generic/unique_by_key.inl b/thrust/thrust/system/detail/generic/unique_by_key.inl
index 2126a9e34fb..cffa5d4df4a 100644
--- a/thrust/thrust/system/detail/generic/unique_by_key.inl
+++ b/thrust/thrust/system/detail/generic/unique_by_key.inl
@@ -51,7 +51,7 @@ _CCCL_HOST_DEVICE thrust::pair<ForwardIterator1, ForwardIterator2> unique_by_key
   ForwardIterator1 keys_last,
   ForwardIterator2 values_first)
 {
-  using KeyType = typename thrust::iterator_traits<ForwardIterator1>::value_type;
+  using KeyType = thrust::detail::it_value_t<ForwardIterator1>;
   return thrust::unique_by_key(exec, keys_first, keys_last, values_first, thrust::equal_to<KeyType>());
 } // end unique_by_key()
 
@@ -63,8 +63,8 @@ _CCCL_HOST_DEVICE thrust::pair<ForwardIterator1, ForwardIterator2> unique_by_key
   ForwardIterator2 values_first,
   BinaryPredicate binary_pred)
 {
-  using InputType1 = typename thrust::iterator_traits<ForwardIterator1>::value_type;
-  using InputType2 = typename thrust::iterator_traits<ForwardIterator2>::value_type;
+  using InputType1 = thrust::detail::it_value_t<ForwardIterator1>;
+  using InputType2 = thrust::detail::it_value_t<ForwardIterator2>;
 
   ForwardIterator2 values_last = values_first + (keys_last - keys_first);
 
@@ -87,7 +87,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> unique_by_key_c
   OutputIterator1 keys_output,
   OutputIterator2 values_output)
 {
-  using KeyType = typename thrust::iterator_traits<InputIterator1>::value_type;
+  using KeyType = thrust::detail::it_value_t<InputIterator1>;
   return thrust::unique_by_key_copy(
     exec, keys_first, keys_last, values_first, keys_output, values_output, thrust::equal_to<KeyType>());
 } // end unique_by_key_copy()
@@ -107,7 +107,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> unique_by_key_c
   OutputIterator2 values_output,
   BinaryPredicate binary_pred)
 {
-  using difference_type = typename thrust::iterator_traits<InputIterator1>::difference_type;
+  using difference_type = thrust::detail::it_difference_t<InputIterator1>;
 
   difference_type n = thrust::distance(keys_first, keys_last);
 
diff --git a/thrust/thrust/system/detail/sequential/adjacent_difference.h b/thrust/thrust/system/detail/sequential/adjacent_difference.h
index f72b390bc32..158293884bf 100644
--- a/thrust/thrust/system/detail/sequential/adjacent_difference.h
+++ b/thrust/thrust/system/detail/sequential/adjacent_difference.h
@@ -49,7 +49,7 @@ _CCCL_HOST_DEVICE OutputIterator adjacent_difference(
   OutputIterator result,
   BinaryFunction binary_op)
 {
-  using InputType = typename thrust::iterator_traits<InputIterator>::value_type;
+  using InputType = thrust::detail::it_value_t<InputIterator>;
 
   if (first == last)
   {
diff --git a/thrust/thrust/system/detail/sequential/binary_search.h b/thrust/thrust/system/detail/sequential/binary_search.h
index b81843ecead..3f390063d3d 100644
--- a/thrust/thrust/system/detail/sequential/binary_search.h
+++ b/thrust/thrust/system/detail/sequential/binary_search.h
@@ -55,7 +55,7 @@ _CCCL_HOST_DEVICE ForwardIterator lower_bound(
   // wrap comp
   thrust::detail::wrapped_function<StrictWeakOrdering, bool> wrapped_comp{comp};
 
-  using difference_type = typename thrust::iterator_difference<ForwardIterator>::type;
+  using difference_type = thrust::detail::it_difference_t<ForwardIterator>;
 
   difference_type len = thrust::distance(first, last);
 
@@ -93,7 +93,7 @@ _CCCL_HOST_DEVICE ForwardIterator upper_bound(
   // wrap comp
   thrust::detail::wrapped_function<StrictWeakOrdering, bool> wrapped_comp{comp};
 
-  using difference_type = typename thrust::iterator_difference<ForwardIterator>::type;
+  using difference_type = thrust::detail::it_difference_t<ForwardIterator>;
 
   difference_type len = thrust::distance(first, last);
 
diff --git a/thrust/thrust/system/detail/sequential/copy.inl b/thrust/thrust/system/detail/sequential/copy.inl
index b6a4958cacf..806088cfac3 100644
--- a/thrust/thrust/system/detail/sequential/copy.inl
+++ b/thrust/thrust/system/detail/sequential/copy.inl
@@ -58,7 +58,7 @@ copy(InputIterator first,
      OutputIterator result,
      thrust::detail::true_type) // is_indirectly_trivially_relocatable_to
 {
-  using Size = typename thrust::iterator_difference<InputIterator>::type;
+  using Size = thrust::detail::it_difference_t<InputIterator>;
 
   const Size n = last - first;
   thrust::system::detail::sequential::trivial_copy_n(get(&*first), n, get(&*result));
diff --git a/thrust/thrust/system/detail/sequential/general_copy.h b/thrust/thrust/system/detail/sequential/general_copy.h
index 990085dd2f4..ffc6588cad5 100644
--- a/thrust/thrust/system/detail/sequential/general_copy.h
+++ b/thrust/thrust/system/detail/sequential/general_copy.h
@@ -42,19 +42,27 @@ namespace sequential
 namespace general_copy_detail
 {
 
-template <typename T1, typename T2>
-struct lazy_is_assignable : ::cuda::std::is_assignable<typename T1::type, typename T2::type>
-{};
-
 // sometimes OutputIterator's reference type is reported as void
 // in that case, just assume that we're able to assign to it OK
 template <typename InputIterator, typename OutputIterator>
 struct reference_is_assignable
-    : thrust::detail::eval_if<
-        ::cuda::std::is_same<typename thrust::iterator_reference<OutputIterator>::type, void>::value,
-        thrust::detail::true_type,
-        lazy_is_assignable<thrust::iterator_reference<OutputIterator>, thrust::iterator_reference<InputIterator>>>::type
-{};
+{
+  template <typename OI>
+  static constexpr bool h()
+  {
+    if constexpr (::cuda::std::is_same_v<thrust::detail::it_reference_t<OI>, void>)
+    {
+      return true;
+    }
+    else
+    {
+      return ::cuda::std::is_assignable_v<thrust::detail::it_reference_t<OI>,
+                                          thrust::detail::it_reference_t<InputIterator>>;
+    }
+  }
+
+  static constexpr bool value = h<OutputIterator>();
+};
 
 // introduce an iterator assign helper to deal with assignments from
 // a wrapped reference
@@ -73,7 +81,7 @@ inline _CCCL_HOST_DEVICE
 typename thrust::detail::disable_if<reference_is_assignable<InputIterator, OutputIterator>::value>::type
 iter_assign(OutputIterator dst, InputIterator src)
 {
-  using value_type = typename thrust::iterator_value<InputIterator>::type;
+  using value_type = thrust::detail::it_value_t<InputIterator>;
 
   // insert a temporary and hope for the best
   *dst = static_cast<value_type>(*src);
diff --git a/thrust/thrust/system/detail/sequential/get_value.h b/thrust/thrust/system/detail/sequential/get_value.h
index 029982f95fd..8444c4f1c61 100644
--- a/thrust/thrust/system/detail/sequential/get_value.h
+++ b/thrust/thrust/system/detail/sequential/get_value.h
@@ -37,7 +37,7 @@ namespace sequential
 {
 
 template <typename DerivedPolicy, typename Pointer>
-_CCCL_HOST_DEVICE typename thrust::iterator_value<Pointer>::type
+_CCCL_HOST_DEVICE thrust::detail::it_value_t<Pointer>
 get_value(sequential::execution_policy<DerivedPolicy>&, Pointer ptr)
 {
   return *thrust::raw_pointer_cast(ptr);
diff --git a/thrust/thrust/system/detail/sequential/insertion_sort.h b/thrust/thrust/system/detail/sequential/insertion_sort.h
index e66a54fd88f..d8e0df86b27 100644
--- a/thrust/thrust/system/detail/sequential/insertion_sort.h
+++ b/thrust/thrust/system/detail/sequential/insertion_sort.h
@@ -42,7 +42,7 @@ _CCCL_EXEC_CHECK_DISABLE
 template <typename RandomAccessIterator, typename StrictWeakOrdering>
 _CCCL_HOST_DEVICE void insertion_sort(RandomAccessIterator first, RandomAccessIterator last, StrictWeakOrdering comp)
 {
-  using value_type = typename thrust::iterator_value<RandomAccessIterator>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator>;
 
   if (first == last)
   {
@@ -86,8 +86,8 @@ template <typename RandomAccessIterator1, typename RandomAccessIterator2, typena
 _CCCL_HOST_DEVICE void insertion_sort_by_key(
   RandomAccessIterator1 first1, RandomAccessIterator1 last1, RandomAccessIterator2 first2, StrictWeakOrdering comp)
 {
-  using value_type1 = typename thrust::iterator_value<RandomAccessIterator1>::type;
-  using value_type2 = typename thrust::iterator_value<RandomAccessIterator2>::type;
+  using value_type1 = thrust::detail::it_value_t<RandomAccessIterator1>;
+  using value_type2 = thrust::detail::it_value_t<RandomAccessIterator2>;
 
   if (first1 == last1)
   {
diff --git a/thrust/thrust/system/detail/sequential/partition.h b/thrust/thrust/system/detail/sequential/partition.h
index cc744bb8866..4f65c6507ae 100644
--- a/thrust/thrust/system/detail/sequential/partition.h
+++ b/thrust/thrust/system/detail/sequential/partition.h
@@ -58,7 +58,7 @@ _CCCL_HOST_DEVICE void iter_swap(ForwardIterator1 iter1, ForwardIterator2 iter2)
 {
   // note: we cannot use swap(*iter1, *iter2) here, because the reference_type's could be proxy references, for which
   // swap() is not guaranteed to work
-  using T = typename thrust::iterator_value<ForwardIterator1>::type;
+  using T = thrust::detail::it_value_t<ForwardIterator1>;
   T temp  = ::cuda::std::move(*iter1);
   *iter1  = ::cuda::std::move(*iter2);
   *iter2  = ::cuda::std::move(temp);
@@ -159,7 +159,7 @@ _CCCL_HOST_DEVICE ForwardIterator stable_partition(
   // wrap pred
   thrust::detail::wrapped_function<Predicate, bool> wrapped_pred{pred};
 
-  using T = typename thrust::iterator_value<ForwardIterator>::type;
+  using T = thrust::detail::it_value_t<ForwardIterator>;
 
   using TempRange    = thrust::detail::temporary_array<T, DerivedPolicy>;
   using TempIterator = typename TempRange::iterator;
@@ -201,7 +201,7 @@ _CCCL_HOST_DEVICE ForwardIterator stable_partition(
   // wrap pred
   thrust::detail::wrapped_function<Predicate, bool> wrapped_pred{pred};
 
-  using T = typename thrust::iterator_value<ForwardIterator>::type;
+  using T = thrust::detail::it_value_t<ForwardIterator>;
 
   using TempRange    = thrust::detail::temporary_array<T, DerivedPolicy>;
   using TempIterator = typename TempRange::iterator;
diff --git a/thrust/thrust/system/detail/sequential/reduce_by_key.h b/thrust/thrust/system/detail/sequential/reduce_by_key.h
index 2318c015cff..bcbc27b2678 100644
--- a/thrust/thrust/system/detail/sequential/reduce_by_key.h
+++ b/thrust/thrust/system/detail/sequential/reduce_by_key.h
@@ -55,11 +55,11 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> reduce_by_key(
   BinaryPredicate binary_pred,
   BinaryFunction binary_op)
 {
-  using InputKeyType   = typename thrust::iterator_traits<InputIterator1>::value_type;
-  using InputValueType = typename thrust::iterator_traits<InputIterator2>::value_type;
+  using InputKeyType   = thrust::detail::it_value_t<InputIterator1>;
+  using InputValueType = thrust::detail::it_value_t<InputIterator2>;
 
   // Use the input iterator's value type per https://wg21.link/P0571
-  using TemporaryType = typename thrust::iterator_value<InputIterator2>::type;
+  using TemporaryType = thrust::detail::it_value_t<InputIterator2>;
 
   if (keys_first != keys_last)
   {
diff --git a/thrust/thrust/system/detail/sequential/scan.h b/thrust/thrust/system/detail/sequential/scan.h
index 86a132ba3e7..599d0e18189 100644
--- a/thrust/thrust/system/detail/sequential/scan.h
+++ b/thrust/thrust/system/detail/sequential/scan.h
@@ -31,7 +31,6 @@
 #endif // no system header
 #include <thrust/detail/function.h>
 #include <thrust/detail/type_traits.h>
-#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/detail/sequential/execution_policy.h>
 
@@ -57,7 +56,7 @@ _CCCL_HOST_DEVICE OutputIterator inclusive_scan(
   using namespace thrust::detail;
 
   // Use the input iterator's value type per https://wg21.link/P0571
-  using ValueType = typename thrust::iterator_value<InputIterator>::type;
+  using ValueType = thrust::detail::it_value_t<InputIterator>;
 
   // wrap binary_op
   thrust::detail::wrapped_function<BinaryFunction, ValueType> wrapped_binary_op{binary_op};
@@ -93,8 +92,8 @@ _CCCL_HOST_DEVICE OutputIterator inclusive_scan(
 {
   using namespace thrust::detail;
 
-  using ValueType = typename ::cuda::std::
-    __accumulator_t<BinaryFunction, typename ::cuda::std::iterator_traits<InputIterator>::value_type, InitialValueType>;
+  using ValueType =
+    typename ::cuda::std::__accumulator_t<BinaryFunction, thrust::detail::it_value_t<InputIterator>, InitialValueType>;
 
   // wrap binary_op
   thrust::detail::wrapped_function<BinaryFunction, ValueType> wrapped_binary_op{binary_op};
diff --git a/thrust/thrust/system/detail/sequential/scan_by_key.h b/thrust/thrust/system/detail/sequential/scan_by_key.h
index 0bb32844091..2d558edb694 100644
--- a/thrust/thrust/system/detail/sequential/scan_by_key.h
+++ b/thrust/thrust/system/detail/sequential/scan_by_key.h
@@ -57,8 +57,8 @@ _CCCL_HOST_DEVICE OutputIterator inclusive_scan_by_key(
   BinaryPredicate binary_pred,
   BinaryFunction binary_op)
 {
-  using KeyType   = typename thrust::iterator_traits<InputIterator1>::value_type;
-  using ValueType = typename thrust::iterator_traits<InputIterator2>::value_type;
+  using KeyType   = thrust::detail::it_value_t<InputIterator1>;
+  using ValueType = thrust::detail::it_value_t<InputIterator2>;
 
   // wrap binary_op
   thrust::detail::wrapped_function<BinaryFunction, ValueType> wrapped_binary_op{binary_op};
@@ -108,7 +108,7 @@ _CCCL_HOST_DEVICE OutputIterator exclusive_scan_by_key(
   BinaryPredicate binary_pred,
   BinaryFunction binary_op)
 {
-  using KeyType   = typename thrust::iterator_traits<InputIterator1>::value_type;
+  using KeyType   = thrust::detail::it_value_t<InputIterator1>;
   using ValueType = T;
 
   if (first1 != last1)
diff --git a/thrust/thrust/system/detail/sequential/sort.inl b/thrust/thrust/system/detail/sequential/sort.inl
index 7c74b33b891..ab107f933d4 100644
--- a/thrust/thrust/system/detail/sequential/sort.inl
+++ b/thrust/thrust/system/detail/sequential/sort.inl
@@ -63,7 +63,7 @@ _CCCL_HOST_DEVICE void stable_sort(
   thrust::system::detail::sequential::stable_primitive_sort(exec, first, last);
 
   // if comp is greater<T> then reverse the keys
-  using KeyType = typename thrust::iterator_traits<RandomAccessIterator>::value_type;
+  using KeyType = thrust::detail::it_value_t<RandomAccessIterator>;
 
   if (needs_reverse<KeyType, StrictWeakOrdering>::value)
   {
@@ -84,7 +84,7 @@ _CCCL_HOST_DEVICE void stable_sort_by_key(
   thrust::detail::true_type)
 {
   // if comp is greater<T> then reverse the keys and values
-  using KeyType = typename thrust::iterator_traits<RandomAccessIterator1>::value_type;
+  using KeyType = thrust::detail::it_value_t<RandomAccessIterator1>;
 
   // note, we also have to reverse the (unordered) input to preserve stability
   if (needs_reverse<KeyType, StrictWeakOrdering>::value)
@@ -151,7 +151,7 @@ _CCCL_HOST_DEVICE void stable_sort(
   // the compilation time of stable_primitive_sort is too expensive to use within a single CUDA thread
   NV_IF_TARGET(
     NV_IS_HOST,
-    (using KeyType = thrust::iterator_value_t<RandomAccessIterator>;
+    (using KeyType = thrust::detail::it_value_t<RandomAccessIterator>;
      sort_detail::use_primitive_sort<KeyType, StrictWeakOrdering> use_primitive_sort;
      sort_detail::stable_sort(exec, first, last, comp, use_primitive_sort);),
     ( // NV_IS_DEVICE:
@@ -173,7 +173,7 @@ _CCCL_HOST_DEVICE void stable_sort_by_key(
   // the compilation time of stable_primitive_sort_by_key is too expensive to use within a single CUDA thread
   NV_IF_TARGET(
     NV_IS_HOST,
-    (using KeyType = thrust::iterator_value_t<RandomAccessIterator1>;
+    (using KeyType = thrust::detail::it_value_t<RandomAccessIterator1>;
      sort_detail::use_primitive_sort<KeyType, StrictWeakOrdering> use_primitive_sort;
      sort_detail::stable_sort_by_key(exec, first1, last1, first2, comp, use_primitive_sort);),
     ( // NV_IS_DEVICE:
diff --git a/thrust/thrust/system/detail/sequential/stable_merge_sort.inl b/thrust/thrust/system/detail/sequential/stable_merge_sort.inl
index e987540c0e2..a35338c3936 100644
--- a/thrust/thrust/system/detail/sequential/stable_merge_sort.inl
+++ b/thrust/thrust/system/detail/sequential/stable_merge_sort.inl
@@ -32,6 +32,7 @@
 #include <thrust/system/detail/sequential/insertion_sort.h>
 
 #include <cuda/std/__algorithm/min.h>
+#include <cuda/std/iterator>
 
 #include <nv/target>
 
@@ -53,7 +54,7 @@ _CCCL_HOST_DEVICE void inplace_merge(
   RandomAccessIterator last,
   StrictWeakOrdering comp)
 {
-  using value_type = typename thrust::iterator_value<RandomAccessIterator>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator>;
 
   thrust::detail::temporary_array<value_type, DerivedPolicy> a(exec, first, middle);
   thrust::detail::temporary_array<value_type, DerivedPolicy> b(exec, middle, last);
@@ -73,8 +74,8 @@ _CCCL_HOST_DEVICE void inplace_merge_by_key(
   RandomAccessIterator2 first2,
   StrictWeakOrdering comp)
 {
-  using value_type1 = typename thrust::iterator_value<RandomAccessIterator1>::type;
-  using value_type2 = typename thrust::iterator_value<RandomAccessIterator2>::type;
+  using value_type1 = thrust::detail::it_value_t<RandomAccessIterator1>;
+  using value_type2 = thrust::detail::it_value_t<RandomAccessIterator2>;
 
   RandomAccessIterator2 middle2 = first2 + (middle1 - first1);
   RandomAccessIterator2 last2   = first2 + (last1 - first1);
@@ -192,8 +193,8 @@ _CCCL_HOST_DEVICE void iterative_stable_merge_sort(
   RandomAccessIterator last,
   StrictWeakOrdering comp)
 {
-  using value_type      = typename thrust::iterator_value<RandomAccessIterator>::type;
-  using difference_type = typename thrust::iterator_difference<RandomAccessIterator>::type;
+  using value_type      = thrust::detail::it_value_t<RandomAccessIterator>;
+  using difference_type = thrust::detail::it_difference_t<RandomAccessIterator>;
 
   difference_type n = last - first;
 
@@ -236,9 +237,9 @@ _CCCL_HOST_DEVICE void iterative_stable_merge_sort_by_key(
   RandomAccessIterator2 values_first,
   StrictWeakOrdering comp)
 {
-  using value_type1     = typename thrust::iterator_value<RandomAccessIterator1>::type;
-  using value_type2     = typename thrust::iterator_value<RandomAccessIterator2>::type;
-  using difference_type = typename thrust::iterator_difference<RandomAccessIterator1>::type;
+  using value_type1     = thrust::detail::it_value_t<RandomAccessIterator1>;
+  using value_type2     = thrust::detail::it_value_t<RandomAccessIterator2>;
+  using difference_type = thrust::detail::it_difference_t<RandomAccessIterator1>;
 
   difference_type n = keys_last - keys_first;
 
diff --git a/thrust/thrust/system/detail/sequential/stable_primitive_sort.inl b/thrust/thrust/system/detail/sequential/stable_primitive_sort.inl
index 333e25f665d..6199fe87598 100644
--- a/thrust/thrust/system/detail/sequential/stable_primitive_sort.inl
+++ b/thrust/thrust/system/detail/sequential/stable_primitive_sort.inl
@@ -44,12 +44,12 @@ namespace stable_primitive_sort_detail
 
 template <typename Iterator>
 struct enable_if_bool_sort
-    : ::cuda::std::enable_if<::cuda::std::is_same<bool, typename thrust::iterator_value<Iterator>::type>::value>
+    : ::cuda::std::enable_if<::cuda::std::is_same<bool, thrust::detail::it_value_t<Iterator>>::value>
 {};
 
 template <typename Iterator>
 struct disable_if_bool_sort
-    : thrust::detail::disable_if<::cuda::std::is_same<bool, typename thrust::iterator_value<Iterator>::type>::value>
+    : thrust::detail::disable_if<::cuda::std::is_same<bool, thrust::detail::it_value_t<Iterator>>::value>
 {};
 
 template <typename DerivedPolicy, typename RandomAccessIterator>
diff --git a/thrust/thrust/system/detail/sequential/stable_radix_sort.inl b/thrust/thrust/system/detail/sequential/stable_radix_sort.inl
index c72a9fab864..62363565b6e 100644
--- a/thrust/thrust/system/detail/sequential/stable_radix_sort.inl
+++ b/thrust/thrust/system/detail/sequential/stable_radix_sort.inl
@@ -193,7 +193,7 @@ inline _CCCL_HOST_DEVICE void radix_shuffle_n(
   Integer bit_shift,
   size_t* histogram)
 {
-  using KeyType = typename thrust::iterator_value<RandomAccessIterator1>::type;
+  using KeyType = thrust::detail::it_value_t<RandomAccessIterator1>;
 
   // note that we are going to mutate the histogram during this sequential scatter
   thrust::scatter(
@@ -221,7 +221,7 @@ _CCCL_HOST_DEVICE void radix_shuffle_n(
   Integer bit_shift,
   size_t* histogram)
 {
-  using KeyType = typename thrust::iterator_value<RandomAccessIterator1>::type;
+  using KeyType = thrust::detail::it_value_t<RandomAccessIterator1>;
 
   // note that we are going to mutate the histogram during this sequential scatter
   thrust::scatter(
@@ -247,7 +247,7 @@ _CCCL_HOST_DEVICE void radix_sort(
   RandomAccessIterator4 vals2,
   const size_t N)
 {
-  using KeyType = typename thrust::iterator_value<RandomAccessIterator1>::type;
+  using KeyType = thrust::detail::it_value_t<RandomAccessIterator1>;
 
   using Encoder     = RadixEncoder<KeyType>;
   using EncodedType = decltype(::cuda::std::declval<Encoder>()(::cuda::std::declval<KeyType>()));
@@ -533,7 +533,7 @@ _CCCL_HOST_DEVICE void radix_sort(
   RandomAccessIterator2 keys2,
   const size_t N)
 {
-  using KeyType = typename thrust::iterator_value<RandomAccessIterator1>::type;
+  using KeyType = thrust::detail::it_value_t<RandomAccessIterator1>;
   radix_sort_dispatcher<sizeof(KeyType)>()(exec, keys1, keys2, N);
 }
 
@@ -550,7 +550,7 @@ _CCCL_HOST_DEVICE void radix_sort(
   RandomAccessIterator4 vals2,
   const size_t N)
 {
-  using KeyType = typename thrust::iterator_value<RandomAccessIterator1>::type;
+  using KeyType = thrust::detail::it_value_t<RandomAccessIterator1>;
   radix_sort_dispatcher<sizeof(KeyType)>()(exec, keys1, keys2, vals1, vals2, N);
 }
 
@@ -560,7 +560,7 @@ template <typename DerivedPolicy, typename RandomAccessIterator>
 _CCCL_HOST_DEVICE void stable_radix_sort(
   sequential::execution_policy<DerivedPolicy>& exec, RandomAccessIterator first, RandomAccessIterator last)
 {
-  using KeyType = typename thrust::iterator_value<RandomAccessIterator>::type;
+  using KeyType = thrust::detail::it_value_t<RandomAccessIterator>;
 
   size_t N = last - first;
 
@@ -576,8 +576,8 @@ _CCCL_HOST_DEVICE void stable_radix_sort_by_key(
   RandomAccessIterator1 last1,
   RandomAccessIterator2 first2)
 {
-  using KeyType   = typename thrust::iterator_value<RandomAccessIterator1>::type;
-  using ValueType = typename thrust::iterator_value<RandomAccessIterator2>::type;
+  using KeyType   = thrust::detail::it_value_t<RandomAccessIterator1>;
+  using ValueType = thrust::detail::it_value_t<RandomAccessIterator2>;
 
   size_t N = last1 - first1;
 
diff --git a/thrust/thrust/system/detail/sequential/unique.h b/thrust/thrust/system/detail/sequential/unique.h
index 6947de13df0..23cd2142fe8 100644
--- a/thrust/thrust/system/detail/sequential/unique.h
+++ b/thrust/thrust/system/detail/sequential/unique.h
@@ -50,7 +50,7 @@ _CCCL_HOST_DEVICE OutputIterator unique_copy(
   OutputIterator output,
   BinaryPredicate binary_pred)
 {
-  using T = typename thrust::iterator_traits<InputIterator>::value_type;
+  using T = thrust::detail::it_value_t<InputIterator>;
 
   if (first != last)
   {
@@ -89,11 +89,11 @@ _CCCL_HOST_DEVICE ForwardIterator unique(
 } // end unique()
 
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator> unique_count(
   sequential::execution_policy<DerivedPolicy>&, ForwardIterator first, ForwardIterator last, BinaryPredicate binary_pred)
 {
-  using T = typename thrust::iterator_traits<ForwardIterator>::value_type;
-  typename thrust::iterator_traits<ForwardIterator>::difference_type count{};
+  using T = thrust::detail::it_value_t<ForwardIterator>;
+  thrust::detail::it_difference_t<ForwardIterator> count{};
 
   if (first != last)
   {
diff --git a/thrust/thrust/system/detail/sequential/unique_by_key.h b/thrust/thrust/system/detail/sequential/unique_by_key.h
index b1213d09799..e9ef53997ae 100644
--- a/thrust/thrust/system/detail/sequential/unique_by_key.h
+++ b/thrust/thrust/system/detail/sequential/unique_by_key.h
@@ -57,8 +57,8 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> unique_by_key_c
   OutputIterator2 values_output,
   BinaryPredicate binary_pred)
 {
-  using InputKeyType    = typename thrust::iterator_traits<InputIterator1>::value_type;
-  using OutputValueType = typename thrust::iterator_traits<OutputIterator2>::value_type;
+  using InputKeyType    = thrust::detail::it_value_t<InputIterator1>;
+  using OutputValueType = thrust::detail::it_value_t<OutputIterator2>;
 
   if (keys_first != keys_last)
   {
diff --git a/thrust/thrust/system/omp/detail/for_each.inl b/thrust/thrust/system/omp/detail/for_each.inl
index 5fe95310c3f..0ba9cd9922a 100644
--- a/thrust/thrust/system/omp/detail/for_each.inl
+++ b/thrust/thrust/system/omp/detail/for_each.inl
@@ -62,7 +62,7 @@ RandomAccessIterator for_each_n(execution_policy<DerivedPolicy>&, RandomAccessIt
   thrust::detail::wrapped_function<UnaryFunction, void> wrapped_f{f};
 
   // use a signed type for the iteration variable or suffer the consequences of warnings
-  using DifferenceType    = typename thrust::iterator_difference<RandomAccessIterator>::type;
+  using DifferenceType    = thrust::detail::it_difference_t<RandomAccessIterator>;
   DifferenceType signed_n = n;
 
   THRUST_PRAGMA_OMP(parallel for)
diff --git a/thrust/thrust/system/omp/detail/reduce.inl b/thrust/thrust/system/omp/detail/reduce.inl
index 06a2b666e31..a899e58b1bb 100644
--- a/thrust/thrust/system/omp/detail/reduce.inl
+++ b/thrust/thrust/system/omp/detail/reduce.inl
@@ -47,7 +47,7 @@ OutputType reduce(execution_policy<DerivedPolicy>& exec,
                   OutputType init,
                   BinaryFunction binary_op)
 {
-  using difference_type = typename thrust::iterator_difference<InputIterator>::type;
+  using difference_type = thrust::detail::it_difference_t<InputIterator>;
 
   const difference_type n = thrust::distance(first, last);
 
diff --git a/thrust/thrust/system/omp/detail/reduce_intervals.inl b/thrust/thrust/system/omp/detail/reduce_intervals.inl
index 44be63ad844..b795d43c86e 100644
--- a/thrust/thrust/system/omp/detail/reduce_intervals.inl
+++ b/thrust/thrust/system/omp/detail/reduce_intervals.inl
@@ -64,7 +64,7 @@ void reduce_intervals(
     "OpenMP compiler support is not enabled");
 
 #if (THRUST_DEVICE_COMPILER_IS_OMP_CAPABLE == THRUST_TRUE)
-  using OutputType = typename thrust::iterator_value<OutputIterator>::type;
+  using OutputType = thrust::detail::it_value_t<OutputIterator>;
 
   // wrap binary_op
   thrust::detail::wrapped_function<BinaryFunction, OutputType> wrapped_binary_op{binary_op};
diff --git a/thrust/thrust/system/omp/detail/sort.inl b/thrust/thrust/system/omp/detail/sort.inl
index 20f4abb07e4..ffdbf2bf16e 100644
--- a/thrust/thrust/system/omp/detail/sort.inl
+++ b/thrust/thrust/system/omp/detail/sort.inl
@@ -56,7 +56,7 @@ void inplace_merge(execution_policy<DerivedPolicy>& exec,
                    RandomAccessIterator last,
                    StrictWeakOrdering comp)
 {
-  using value_type = typename thrust::iterator_value<RandomAccessIterator>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator>;
 
   thrust::detail::temporary_array<value_type, DerivedPolicy> a(exec, first, middle);
   thrust::detail::temporary_array<value_type, DerivedPolicy> b(exec, middle, last);
@@ -76,8 +76,8 @@ void inplace_merge_by_key(
   RandomAccessIterator2 first2,
   StrictWeakOrdering comp)
 {
-  using value_type1 = typename thrust::iterator_value<RandomAccessIterator1>::type;
-  using value_type2 = typename thrust::iterator_value<RandomAccessIterator2>::type;
+  using value_type1 = thrust::detail::it_value_t<RandomAccessIterator1>;
+  using value_type2 = thrust::detail::it_value_t<RandomAccessIterator2>;
 
   RandomAccessIterator2 middle2 = first2 + (middle1 - first1);
   RandomAccessIterator2 last2   = first2 + (last1 - first1);
@@ -109,7 +109,7 @@ void stable_sort(
 
   // Avoid issues on compilers that don't provide `omp_get_num_threads()`.
 #if (THRUST_DEVICE_COMPILER_IS_OMP_CAPABLE == THRUST_TRUE)
-  using IndexType = typename thrust::iterator_difference<RandomAccessIterator>::type;
+  using IndexType = thrust::detail::it_difference_t<RandomAccessIterator>;
 
   if (first == last)
   {
@@ -188,7 +188,7 @@ void stable_sort_by_key(
 
   // Avoid issues on compilers that don't provide `omp_get_num_threads()`.
 #if (THRUST_DEVICE_COMPILER_IS_OMP_CAPABLE == THRUST_TRUE)
-  using IndexType = typename thrust::iterator_difference<RandomAccessIterator1>::type;
+  using IndexType = thrust::detail::it_difference_t<RandomAccessIterator1>;
 
   if (keys_first == keys_last)
   {
diff --git a/thrust/thrust/system/omp/detail/unique.h b/thrust/thrust/system/omp/detail/unique.h
index c3cddb418db..1c70e789762 100644
--- a/thrust/thrust/system/omp/detail/unique.h
+++ b/thrust/thrust/system/omp/detail/unique.h
@@ -49,7 +49,7 @@ OutputIterator unique_copy(
   BinaryPredicate binary_pred);
 
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+thrust::detail::it_difference_t<ForwardIterator> unique_count(
   execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last, BinaryPredicate binary_pred);
 
 } // end namespace detail
diff --git a/thrust/thrust/system/omp/detail/unique.inl b/thrust/thrust/system/omp/detail/unique.inl
index 22ac1801a3b..39bf8427f6b 100644
--- a/thrust/thrust/system/omp/detail/unique.inl
+++ b/thrust/thrust/system/omp/detail/unique.inl
@@ -58,7 +58,7 @@ OutputIterator unique_copy(
 } // end unique_copy()
 
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+thrust::detail::it_difference_t<ForwardIterator> unique_count(
   execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last, BinaryPredicate binary_pred)
 {
   // omp prefers generic::unique_count to cpp::unique_count
diff --git a/thrust/thrust/system/tbb/detail/copy_if.inl b/thrust/thrust/system/tbb/detail/copy_if.inl
index 79bf2300b71..1557e6a2b4b 100644
--- a/thrust/thrust/system/tbb/detail/copy_if.inl
+++ b/thrust/thrust/system/tbb/detail/copy_if.inl
@@ -115,7 +115,7 @@ template <typename InputIterator1, typename InputIterator2, typename OutputItera
 OutputIterator
 copy_if(tag, InputIterator1 first, InputIterator1 last, InputIterator2 stencil, OutputIterator result, Predicate pred)
 {
-  using Size = typename thrust::iterator_difference<InputIterator1>::type;
+  using Size = thrust::detail::it_difference_t<InputIterator1>;
   using Body = typename copy_if_detail::body<InputIterator1, InputIterator2, OutputIterator, Predicate, Size>;
 
   Size n = thrust::distance(first, last);
diff --git a/thrust/thrust/system/tbb/detail/reduce.inl b/thrust/thrust/system/tbb/detail/reduce.inl
index 3a45ab246cb..6241dad5058 100644
--- a/thrust/thrust/system/tbb/detail/reduce.inl
+++ b/thrust/thrust/system/tbb/detail/reduce.inl
@@ -114,7 +114,7 @@ template <typename DerivedPolicy, typename InputIterator, typename OutputType, t
 OutputType reduce(
   execution_policy<DerivedPolicy>&, InputIterator begin, InputIterator end, OutputType init, BinaryFunction binary_op)
 {
-  using Size = typename thrust::iterator_difference<InputIterator>::type;
+  using Size = thrust::detail::it_difference_t<InputIterator>;
 
   Size n = thrust::distance(begin, end);
 
diff --git a/thrust/thrust/system/tbb/detail/reduce_by_key.inl b/thrust/thrust/system/tbb/detail/reduce_by_key.inl
index 96c4630ae8f..100e8e18bbc 100644
--- a/thrust/thrust/system/tbb/detail/reduce_by_key.inl
+++ b/thrust/thrust/system/tbb/detail/reduce_by_key.inl
@@ -28,7 +28,6 @@
 #include <thrust/detail/range/tail_flags.h>
 #include <thrust/detail/seq.h>
 #include <thrust/detail/temporary_array.h>
-#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
 #include <thrust/iterator/reverse_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/system/tbb/detail/execution_policy.h>
@@ -64,12 +63,12 @@ inline L divide_ri(const L x, const R y)
 template <typename InputIterator, typename BinaryFunction, typename SFINAE = void>
 struct partial_sum_type
 {
-  using type = thrust::iterator_value_t<InputIterator>;
+  using type = thrust::detail::it_value_t<InputIterator>;
 };
 
 template <typename InputIterator1, typename InputIterator2, typename BinaryPredicate, typename BinaryFunction>
 thrust::pair<InputIterator1,
-             thrust::pair<thrust::iterator_value_t<InputIterator1>,
+             thrust::pair<thrust::detail::it_value_t<InputIterator1>,
                           typename partial_sum_type<InputIterator2, BinaryFunction>::type>>
 reduce_last_segment_backward(
   InputIterator1 keys_first,
@@ -78,14 +77,14 @@ reduce_last_segment_backward(
   BinaryPredicate binary_pred,
   BinaryFunction binary_op)
 {
-  typename thrust::iterator_difference<InputIterator1>::type n = keys_last - keys_first;
+  thrust::detail::it_difference_t<InputIterator1> n = keys_last - keys_first;
 
   // reverse the ranges and consume from the end
   thrust::reverse_iterator<InputIterator1> keys_first_r(keys_last);
   thrust::reverse_iterator<InputIterator1> keys_last_r(keys_first);
   thrust::reverse_iterator<InputIterator2> values_first_r(values_first + n);
 
-  thrust::iterator_value_t<InputIterator1> result_key                          = *keys_first_r;
+  thrust::detail::it_value_t<InputIterator1> result_key                        = *keys_first_r;
   typename partial_sum_type<InputIterator2, BinaryFunction>::type result_value = *values_first_r;
 
   // consume the entirety of the first key's sequence
@@ -106,7 +105,7 @@ template <typename InputIterator1,
           typename BinaryFunction>
 thrust::tuple<OutputIterator1,
               OutputIterator2,
-              thrust::iterator_value_t<InputIterator1>,
+              thrust::detail::it_value_t<InputIterator1>,
               typename partial_sum_type<InputIterator2, BinaryFunction>::type>
 reduce_by_key_with_carry(
   InputIterator1 keys_first,
@@ -119,7 +118,8 @@ reduce_by_key_with_carry(
 {
   // first, consume the last sequence to produce the carry
   // XXX is there an elegant way to pose this such that we don't need to default construct carry?
-  thrust::pair<thrust::iterator_value_t<InputIterator1>, typename partial_sum_type<InputIterator2, BinaryFunction>::type>
+  thrust::pair<thrust::detail::it_value_t<InputIterator1>,
+               typename partial_sum_type<InputIterator2, BinaryFunction>::type>
     carry;
 
   thrust::tie(keys_last, carry) =
@@ -150,7 +150,7 @@ template <typename Iterator1,
           typename BinaryFunction>
 struct serial_reduce_by_key_body
 {
-  using size_type = typename thrust::iterator_difference<Iterator1>::type;
+  using size_type = thrust::detail::it_difference_t<Iterator1>;
 
   Iterator1 keys_first;
   Iterator2 values_first;
@@ -209,7 +209,7 @@ struct serial_reduce_by_key_body
     Iterator6 my_carry_result  = carry_result + interval_idx;
 
     // consume the rest of the interval with reduce_by_key
-    using key_type   = thrust::iterator_value_t<Iterator1>;
+    using key_type   = thrust::detail::it_value_t<Iterator1>;
     using value_type = typename partial_sum_type<Iterator2, BinaryFunction>::type;
 
     // XXX is there a way to pose this so that we don't require default construction of carry?
@@ -255,7 +255,7 @@ make_serial_reduce_by_key_body(
   Iterator4 keys_result,
   Iterator5 values_result,
   Iterator6 carry_result,
-  typename thrust::iterator_difference<Iterator1>::type n,
+  thrust::detail::it_difference_t<Iterator1> n,
   size_t interval_size,
   size_t num_intervals,
   BinaryPredicate binary_pred,
@@ -302,7 +302,7 @@ thrust::pair<Iterator3, Iterator4> reduce_by_key(
   BinaryPredicate binary_pred,
   BinaryFunction binary_op)
 {
-  using difference_type = typename thrust::iterator_difference<Iterator1>::type;
+  using difference_type = thrust::detail::it_difference_t<Iterator1>;
   difference_type n     = keys_last - keys_first;
   if (n == 0)
   {
diff --git a/thrust/thrust/system/tbb/detail/reduce_intervals.h b/thrust/thrust/system/tbb/detail/reduce_intervals.h
index d29b830823b..e3552aca2fb 100644
--- a/thrust/thrust/system/tbb/detail/reduce_intervals.h
+++ b/thrust/thrust/system/tbb/detail/reduce_intervals.h
@@ -112,7 +112,7 @@ void reduce_intervals(
   RandomAccessIterator2 result,
   BinaryFunction binary_op)
 {
-  typename thrust::iterator_difference<RandomAccessIterator1>::type n = last - first;
+  thrust::detail::it_difference_t<RandomAccessIterator1> n = last - first;
 
   Size num_intervals = reduce_intervals_detail::divide_ri(n, interval_size);
 
@@ -129,7 +129,7 @@ void reduce_intervals(
   Size interval_size,
   RandomAccessIterator2 result)
 {
-  using value_type = typename thrust::iterator_value<RandomAccessIterator1>::type;
+  using value_type = thrust::detail::it_value_t<RandomAccessIterator1>;
 
   return thrust::system::tbb::detail::reduce_intervals(
     exec, first, last, interval_size, result, thrust::plus<value_type>());
diff --git a/thrust/thrust/system/tbb/detail/scan.inl b/thrust/thrust/system/tbb/detail/scan.inl
index 09b3cc1d25d..09cbc3a6eb4 100644
--- a/thrust/thrust/system/tbb/detail/scan.inl
+++ b/thrust/thrust/system/tbb/detail/scan.inl
@@ -28,7 +28,6 @@
 #include <thrust/advance.h>
 #include <thrust/detail/function.h>
 #include <thrust/detail/type_traits.h>
-#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
 #include <thrust/distance.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/tbb/detail/scan.h>
@@ -241,9 +240,9 @@ inclusive_scan(tag, InputIterator first, InputIterator last, OutputIterator resu
   using namespace thrust::detail;
 
   // Use the input iterator's value type per https://wg21.link/P0571
-  using ValueType = typename thrust::iterator_value<InputIterator>::type;
+  using ValueType = thrust::detail::it_value_t<InputIterator>;
 
-  using Size = typename thrust::iterator_difference<InputIterator>::type;
+  using Size = thrust::detail::it_difference_t<InputIterator>;
   Size n     = thrust::distance(first, last);
 
   if (n != 0)
@@ -265,10 +264,10 @@ OutputIterator inclusive_scan(
   using namespace thrust::detail;
 
   // Use the input iterator's value type and the initial value type per wg21.link/p2322
-  using ValueType = typename ::cuda::std::
-    __accumulator_t<BinaryFunction, typename ::cuda::std::iterator_traits<InputIterator>::value_type, InitialValueType>;
+  using ValueType =
+    typename ::cuda::std::__accumulator_t<BinaryFunction, thrust::detail::it_value_t<InputIterator>, InitialValueType>;
 
-  using Size = typename thrust::iterator_difference<InputIterator>::type;
+  using Size = thrust::detail::it_difference_t<InputIterator>;
   Size n     = thrust::distance(first, last);
 
   if (n != 0)
@@ -292,7 +291,7 @@ OutputIterator exclusive_scan(
   // Use the initial value type per https://wg21.link/P0571
   using ValueType = InitialValueType;
 
-  using Size = typename thrust::iterator_difference<InputIterator>::type;
+  using Size = thrust::detail::it_difference_t<InputIterator>;
   Size n     = thrust::distance(first, last);
 
   if (n != 0)
diff --git a/thrust/thrust/system/tbb/detail/sort.inl b/thrust/thrust/system/tbb/detail/sort.inl
index a6485ed2be7..aa7c148713d 100644
--- a/thrust/thrust/system/tbb/detail/sort.inl
+++ b/thrust/thrust/system/tbb/detail/sort.inl
@@ -94,7 +94,7 @@ void merge_sort(execution_policy<DerivedPolicy>& exec,
                 StrictWeakOrdering comp,
                 bool inplace)
 {
-  using difference_type = typename thrust::iterator_difference<Iterator1>::type;
+  using difference_type = thrust::detail::it_difference_t<Iterator1>;
 
   difference_type n = thrust::distance(first1, last1);
 
@@ -212,7 +212,7 @@ void merge_sort_by_key(
   StrictWeakOrdering comp,
   bool inplace)
 {
-  using difference_type = typename thrust::iterator_difference<Iterator1>::type;
+  using difference_type = thrust::detail::it_difference_t<Iterator1>;
 
   difference_type n = thrust::distance(first1, last1);
 
@@ -260,7 +260,7 @@ template <typename DerivedPolicy, typename RandomAccessIterator, typename Strict
 void stable_sort(
   execution_policy<DerivedPolicy>& exec, RandomAccessIterator first, RandomAccessIterator last, StrictWeakOrdering comp)
 {
-  using key_type = typename thrust::iterator_value<RandomAccessIterator>::type;
+  using key_type = thrust::detail::it_value_t<RandomAccessIterator>;
 
   thrust::detail::temporary_array<key_type, DerivedPolicy> temp(exec, first, last);
 
@@ -278,8 +278,8 @@ void stable_sort_by_key(
   RandomAccessIterator2 first2,
   StrictWeakOrdering comp)
 {
-  using key_type = typename thrust::iterator_value<RandomAccessIterator1>::type;
-  using val_type = typename thrust::iterator_value<RandomAccessIterator2>::type;
+  using key_type = thrust::detail::it_value_t<RandomAccessIterator1>;
+  using val_type = thrust::detail::it_value_t<RandomAccessIterator2>;
 
   RandomAccessIterator2 last2 = first2 + thrust::distance(first1, last1);
 
diff --git a/thrust/thrust/system/tbb/detail/unique.h b/thrust/thrust/system/tbb/detail/unique.h
index 4cf5f6cf40b..17730cb7d94 100644
--- a/thrust/thrust/system/tbb/detail/unique.h
+++ b/thrust/thrust/system/tbb/detail/unique.h
@@ -49,7 +49,7 @@ OutputIterator unique_copy(
   BinaryPredicate binary_pred);
 
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+thrust::detail::it_difference_t<ForwardIterator> unique_count(
   execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last, BinaryPredicate binary_pred);
 
 } // end namespace detail
diff --git a/thrust/thrust/system/tbb/detail/unique.inl b/thrust/thrust/system/tbb/detail/unique.inl
index 051263b9edc..0fe6171f3d6 100644
--- a/thrust/thrust/system/tbb/detail/unique.inl
+++ b/thrust/thrust/system/tbb/detail/unique.inl
@@ -58,7 +58,7 @@ OutputIterator unique_copy(
 } // end unique_copy()
 
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+thrust::detail::it_difference_t<ForwardIterator> unique_count(
   execution_policy<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last, BinaryPredicate binary_pred)
 {
   // tbb prefers generic::unique_count to cpp::unique_count
diff --git a/thrust/thrust/type_traits/is_trivially_relocatable.h b/thrust/thrust/type_traits/is_trivially_relocatable.h
index 616cd83805d..d0943b8b368 100644
--- a/thrust/thrust/type_traits/is_trivially_relocatable.h
+++ b/thrust/thrust/type_traits/is_trivially_relocatable.h
@@ -144,11 +144,10 @@ constexpr bool is_trivially_relocatable_to_v = is_trivially_relocatable_to<From,
  * \see THRUST_PROCLAIM_TRIVIALLY_RELOCATABLE
  */
 template <typename FromIterator, typename ToIterator>
-using is_indirectly_trivially_relocatable_to =
-  integral_constant<bool,
-                    is_contiguous_iterator<FromIterator>::value && is_contiguous_iterator<ToIterator>::value
-                      && is_trivially_relocatable_to<typename thrust::iterator_traits<FromIterator>::value_type,
-                                                     typename thrust::iterator_traits<ToIterator>::value_type>::value>;
+using is_indirectly_trivially_relocatable_to = integral_constant<
+  bool,
+  is_contiguous_iterator<FromIterator>::value && is_contiguous_iterator<ToIterator>::value
+    && is_trivially_relocatable_to<detail::it_value_t<FromIterator>, detail::it_value_t<ToIterator>>::value>;
 
 /*! \brief <tt>constexpr bool</tt> that is \c true if the element type of
  *  \c FromIterator is
diff --git a/thrust/thrust/unique.h b/thrust/thrust/unique.h
index 6310b17ebd9..c5721a88fab 100644
--- a/thrust/thrust/unique.h
+++ b/thrust/thrust/unique.h
@@ -961,7 +961,7 @@ thrust::pair<OutputIterator1, OutputIterator2> unique_by_key_copy(
  *  \see reduce_by_key_copy
  */
 template <typename DerivedPolicy, typename ForwardIterator, typename BinaryPredicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator> unique_count(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
   ForwardIterator first,
   ForwardIterator last,
@@ -1005,7 +1005,7 @@ _CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_
  *  \see reduce_by_key_copy
  */
 template <typename DerivedPolicy, typename ForwardIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type unique_count(
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator> unique_count(
   const thrust::detail::execution_policy_base<DerivedPolicy>& exec, ForwardIterator first, ForwardIterator last);
 
 /*! \p unique_count counts runs of equal elements in the range <tt>[first, last)</tt>
@@ -1042,7 +1042,7 @@ _CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_
  *  \see reduce_by_key_copy
  */
 template <typename ForwardIterator, typename BinaryPredicate>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator>
 unique_count(ForwardIterator first, ForwardIterator last, BinaryPredicate binary_pred);
 
 /*! \p unique_count counts runs of equal elements in the range <tt>[first, last)</tt>
@@ -1079,7 +1079,7 @@ unique_count(ForwardIterator first, ForwardIterator last, BinaryPredicate binary
  *  \see reduce_by_key_copy
  */
 template <typename ForwardIterator>
-_CCCL_HOST_DEVICE typename thrust::iterator_traits<ForwardIterator>::difference_type
+_CCCL_HOST_DEVICE thrust::detail::it_difference_t<ForwardIterator>
 unique_count(ForwardIterator first, ForwardIterator last);
 
 /*! \} // end stream_compaction

From c45663a8a2e71d74629282716a0483f4f99ee799 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Tue, 25 Feb 2025 16:05:25 +0100
Subject: [PATCH 2/2] Change examples to use libcu++

---
 thrust/examples/expand.cu         | 2 +-
 thrust/examples/repeated_range.cu | 2 +-
 thrust/examples/strided_range.cu  | 2 +-
 thrust/examples/tiled_range.cu    | 2 +-
 4 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu
index 350ec1d4099..f27c37d68a5 100644
--- a/thrust/examples/expand.cu
+++ b/thrust/examples/expand.cu
@@ -20,7 +20,7 @@
 template <typename InputIterator1, typename InputIterator2, typename OutputIterator>
 OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, OutputIterator output)
 {
-  using difference_type = thrust::detail::it_difference_t<InputIterator1>;
+  using difference_type = typename cuda::std::iterator_traits<InputIterator1>::difference_type;
 
   difference_type input_size  = thrust::distance(first1, last1);
   difference_type output_size = thrust::reduce(first1, last1);
diff --git a/thrust/examples/repeated_range.cu b/thrust/examples/repeated_range.cu
index debd13c58f9..e454d987314 100644
--- a/thrust/examples/repeated_range.cu
+++ b/thrust/examples/repeated_range.cu
@@ -21,7 +21,7 @@ template <typename Iterator>
 class repeated_range
 {
 public:
-  using difference_type = thrust::detail::it_difference_t<Iterator>;
+  using difference_type = typename cuda::std::iterator_traits<Iterator>::difference_type;
 
   struct repeat_functor
   {
diff --git a/thrust/examples/strided_range.cu b/thrust/examples/strided_range.cu
index dbdbac82d62..7389e1f0ae3 100644
--- a/thrust/examples/strided_range.cu
+++ b/thrust/examples/strided_range.cu
@@ -21,7 +21,7 @@ template <typename Iterator>
 class strided_range
 {
 public:
-  using difference_type = thrust::detail::it_difference_t<Iterator>;
+  using difference_type = typename cuda::std::iterator_traits<Iterator>::difference_type;
 
   struct stride_functor
   {
diff --git a/thrust/examples/tiled_range.cu b/thrust/examples/tiled_range.cu
index 220591dc8f5..7f67a409cb4 100644
--- a/thrust/examples/tiled_range.cu
+++ b/thrust/examples/tiled_range.cu
@@ -21,7 +21,7 @@ template <typename Iterator>
 class tiled_range
 {
 public:
-  using difference_type = thrust::detail::it_difference_t<Iterator>;
+  using difference_type = typename cuda::std::iterator_traits<Iterator>::difference_type;
 
   struct tile_functor
   {