From eca53a22ed3e4adc8a433a113a302fa5d823dc32 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Mon, 24 Feb 2025 13:40:48 +0100
Subject: [PATCH 01/17] Minimize usage of cub::Traits (#3863)

* Replace all uses of cub::Traits other than radix sort key twiddling by numeric_limits
* Drop obsolete specializations of cub::NumericTraits
* Fix radix sort custom type example mentioning non-existent cub::RadixTraits
* Replace cub::BaseTraits and cub::Traits by aliases so uses can no longer specialize it
* Deprecate cub::Traits::Max|Lowest
* Extend documentation of trait classes
* Fix: readd template parameter to cub::BaseTraits to specify whether a type is primitive
* Make is_primitive depend on cub::Traits again

Fixes: #920
---
 c2h/include/c2h/bfloat16.cuh                  |   4 +-
 c2h/include/c2h/half.cuh                      |   5 +-
 c2h/include/c2h/test_util_vec.h               |  19 +-
 cub/benchmarks/bench/reduce/arg_extrema.cu    |   5 +-
 cub/cub/device/device_reduce.cuh              |  23 +-
 cub/cub/device/device_segmented_reduce.cuh    |  21 +-
 cub/cub/util_type.cuh                         | 225 +++++++++++-------
 .../example_device_radix_sort_custom.cu       |  12 +-
 cub/test/catch2_test_device_reduce.cuh        |  23 --
 .../catch2_test_device_segmented_sort_keys.cu |   2 +
 cub/test/catch2_test_util_type.cu             |  27 +++
 cub/test/test_util.h                          |  42 ----
 12 files changed, 196 insertions(+), 212 deletions(-)

diff --git a/c2h/include/c2h/bfloat16.cuh b/c2h/include/c2h/bfloat16.cuh
index 6767850b373..de7f57158f6 100644
--- a/c2h/include/c2h/bfloat16.cuh
+++ b/c2h/include/c2h/bfloat16.cuh
@@ -266,12 +266,10 @@ public:
 };
 _LIBCUDACXX_END_NAMESPACE_STD
 
-_CCCL_SUPPRESS_DEPRECATED_PUSH
 template <>
 struct CUB_NS_QUALIFIER::NumericTraits<bfloat16_t>
-    : CUB_NS_QUALIFIER::BaseTraits<FLOATING_POINT, unsigned short, bfloat16_t>
+    : CUB_NS_QUALIFIER::BaseTraits<FLOATING_POINT, true, unsigned short, bfloat16_t>
 {};
-_CCCL_SUPPRESS_DEPRECATED_POP
 
 #ifdef __GNUC__
 #  pragma GCC diagnostic pop
diff --git a/c2h/include/c2h/half.cuh b/c2h/include/c2h/half.cuh
index 4a30202fe3a..b29f3104a84 100644
--- a/c2h/include/c2h/half.cuh
+++ b/c2h/include/c2h/half.cuh
@@ -361,11 +361,10 @@ public:
 };
 _LIBCUDACXX_END_NAMESPACE_STD
 
-_CCCL_SUPPRESS_DEPRECATED_PUSH
 template <>
-struct CUB_NS_QUALIFIER::NumericTraits<half_t> : CUB_NS_QUALIFIER::BaseTraits<FLOATING_POINT, unsigned short, half_t>
+struct CUB_NS_QUALIFIER::NumericTraits<half_t>
+    : CUB_NS_QUALIFIER::BaseTraits<FLOATING_POINT, true, unsigned short, half_t>
 {};
-_CCCL_SUPPRESS_DEPRECATED_POP
 
 #ifdef __GNUC__
 #  pragma GCC diagnostic pop
diff --git a/c2h/include/c2h/test_util_vec.h b/c2h/include/c2h/test_util_vec.h
index 42e5a33ef7e..01022be9777 100644
--- a/c2h/include/c2h/test_util_vec.h
+++ b/c2h/include/c2h/test_util_vec.h
@@ -289,7 +289,7 @@ C2H_VEC_OVERLOAD(ulonglong, unsigned long long)
 C2H_VEC_OVERLOAD(float, float)
 C2H_VEC_OVERLOAD(double, double)
 
-// Specialize cub::NumericTraits and cuda::std::numeric_limits for vector types.
+// Specialize cuda::std::numeric_limits for vector types.
 
 #  define REPEAT_TO_LIST_1(a)  a
 #  define REPEAT_TO_LIST_2(a)  a, a
@@ -298,23 +298,6 @@ C2H_VEC_OVERLOAD(double, double)
 #  define REPEAT_TO_LIST(N, a) _CCCL_PP_CAT(REPEAT_TO_LIST_, N)(a)
 
 #  define C2H_VEC_TRAITS_OVERLOAD_IMPL(T, BaseT, N)                               \
-    CUB_NAMESPACE_BEGIN                                                           \
-    template <>                                                                   \
-    struct NumericTraits<T>                                                       \
-    {                                                                             \
-      static __host__ __device__ T Max()                                          \
-      {                                                                           \
-        T retval = {REPEAT_TO_LIST(N, NumericTraits<BaseT>::Max())};              \
-        return retval;                                                            \
-      }                                                                           \
-      static __host__ __device__ T Lowest()                                       \
-      {                                                                           \
-        T retval = {REPEAT_TO_LIST(N, NumericTraits<BaseT>::Lowest())};           \
-        return retval;                                                            \
-      }                                                                           \
-    };                                                                            \
-    CUB_NAMESPACE_END                                                             \
-                                                                                  \
     _LIBCUDACXX_BEGIN_NAMESPACE_STD                                               \
     template <>                                                                   \
     class numeric_limits<T>                                                       \
diff --git a/cub/benchmarks/bench/reduce/arg_extrema.cu b/cub/benchmarks/bench/reduce/arg_extrema.cu
index 428dc8b2538..37adfe275ab 100644
--- a/cub/benchmarks/bench/reduce/arg_extrema.cu
+++ b/cub/benchmarks/bench/reduce/arg_extrema.cu
@@ -4,6 +4,7 @@
 #include <cub/device/device_reduce.cuh>
 #include <cub/device/dispatch/dispatch_streaming_reduce.cuh>
 
+#include <cuda/std/limits>
 #include <cuda/std/type_traits>
 
 #include <nvbench_helper.cuh>
@@ -57,7 +58,9 @@ struct policy_hub_t
     // Type used for the final result
     using output_tuple_t = cub::KeyValuePair<global_offset_t, T>;
 
-    auto const init = ::cuda::std::is_same_v<OpT, cub::ArgMin> ? cub::Traits<T>::Max() : cub::Traits<T>::Lowest();
+    auto const init = ::cuda::std::is_same_v<OpT, cub::ArgMin>
+                      ? ::cuda::std::numeric_limits<T>::max()
+                      : ::cuda::std::numeric_limits<T>::lowest();
 
 #if !TUNE_BASE
     using policy_t   = policy_hub_t<output_tuple_t, per_partition_offset_t>;
diff --git a/cub/cub/device/device_reduce.cuh b/cub/cub/device/device_reduce.cuh
index 4a57ca3d7f2..93eb2024f67 100644
--- a/cub/cub/device/device_reduce.cuh
+++ b/cub/cub/device/device_reduce.cuh
@@ -51,6 +51,8 @@
 
 #include <thrust/iterator/tabulate_output_iterator.h>
 
+#include <cuda/std/limits>
+
 #include <iterator>
 
 CUB_NAMESPACE_BEGIN
@@ -334,7 +336,7 @@ struct DeviceReduce
   //! @rst
   //! Computes a device-wide minimum using the less-than (``<``) operator.
   //!
-  //! - Uses ``std::numeric_limits<T>::max()`` as the initial value of the reduction.
+  //! - Uses ``::cuda::std::numeric_limits<T>::max()`` as the initial value of the reduction.
   //! - Does not support ``<`` operators that are non-commutative.
   //! - Provides "run-to-run" determinism for pseudo-associative reduction
   //!   (e.g., addition of floating point types) on the same GPU device.
@@ -433,8 +435,7 @@ struct DeviceReduce
       d_out,
       static_cast<OffsetT>(num_items),
       ::cuda::minimum<>{},
-      // TODO(bgruber): replace with ::cuda::std::numeric_limits<T>::max() (breaking change)
-      Traits<InitT>::Max(),
+      ::cuda::std::numeric_limits<InitT>::max(),
       stream);
   }
 
@@ -583,7 +584,7 @@ struct DeviceReduce
   //!   (assuming the value type of ``d_in`` is ``T``)
   //!
   //!   - The minimum is written to ``d_out.value`` and its offset in the input array is written to ``d_out.key``.
-  //!   - The ``{1, std::numeric_limits<T>::max()}`` tuple is produced for zero-length inputs
+  //!   - The ``{1, ::cuda::std::numeric_limits<T>::max()}`` tuple is produced for zero-length inputs
   //!
   //! - Does not support ``<`` operators that are non-commutative.
   //! - Provides "run-to-run" determinism for pseudo-associative reduction
@@ -690,8 +691,7 @@ struct DeviceReduce
     ArgIndexInputIteratorT d_indexed_in(d_in);
 
     // Initial value
-    // TODO Address https://github.com/NVIDIA/cub/issues/651
-    InitT initial_value{AccumT(1, Traits<InputValueT>::Max())};
+    InitT initial_value{AccumT(1, ::cuda::std::numeric_limits<InputValueT>::max())};
 
     return DispatchReduce<ArgIndexInputIteratorT, OutputIteratorT, OffsetT, cub::ArgMin, InitT, AccumT>::Dispatch(
       d_temp_storage, temp_storage_bytes, d_indexed_in, d_out, num_items, cub::ArgMin(), initial_value, stream);
@@ -700,7 +700,7 @@ struct DeviceReduce
   //! @rst
   //! Computes a device-wide maximum using the greater-than (``>``) operator.
   //!
-  //! - Uses ``std::numeric_limits<T>::lowest()`` as the initial value of the reduction.
+  //! - Uses ``::cuda::std::numeric_limits<T>::lowest()`` as the initial value of the reduction.
   //! - Does not support ``>`` operators that are non-commutative.
   //! - Provides "run-to-run" determinism for pseudo-associative reduction
   //!   (e.g., addition of floating point types) on the same GPU device.
@@ -796,8 +796,7 @@ struct DeviceReduce
       d_out,
       static_cast<OffsetT>(num_items),
       ::cuda::maximum<>{},
-      // TODO(bgruber): replace with ::cuda::std::numeric_limits<T>::lowest() (breaking change)
-      Traits<InitT>::Lowest(),
+      ::cuda::std::numeric_limits<InitT>::lowest(),
       stream);
   }
 
@@ -948,7 +947,7 @@ struct DeviceReduce
   //!
   //!   - The maximum is written to ``d_out.value`` and its offset in the input
   //!     array is written to ``d_out.key``.
-  //!   - The ``{1, std::numeric_limits<T>::lowest()}`` tuple is produced for zero-length inputs
+  //!   - The ``{1, ::cuda::std::numeric_limits<T>::lowest()}`` tuple is produced for zero-length inputs
   //!
   //! - Does not support ``>`` operators that are non-commutative.
   //! - Provides "run-to-run" determinism for pseudo-associative reduction
@@ -1057,9 +1056,7 @@ struct DeviceReduce
     ArgIndexInputIteratorT d_indexed_in(d_in);
 
     // Initial value
-    // TODO Address https://github.com/NVIDIA/cub/issues/651
-    // TODO(bgruber): replace with ::cuda::std::numeric_limits<T>::lowest() (breaking change)
-    InitT initial_value{AccumT(1, Traits<InputValueT>::Lowest())};
+    InitT initial_value{AccumT(1, ::cuda::std::numeric_limits<InputValueT>::lowest())};
 
     return DispatchReduce<ArgIndexInputIteratorT, OutputIteratorT, OffsetT, cub::ArgMax, InitT, AccumT>::Dispatch(
       d_temp_storage, temp_storage_bytes, d_indexed_in, d_out, num_items, cub::ArgMax(), initial_value, stream);
diff --git a/cub/cub/device/device_segmented_reduce.cuh b/cub/cub/device/device_segmented_reduce.cuh
index 642f2521d97..551c487a847 100644
--- a/cub/cub/device/device_segmented_reduce.cuh
+++ b/cub/cub/device/device_segmented_reduce.cuh
@@ -49,6 +49,7 @@
 #include <cub/iterator/arg_index_input_iterator.cuh>
 #include <cub/util_type.cuh>
 
+#include <cuda/std/limits>
 #include <cuda/std/type_traits>
 
 #include <iterator>
@@ -392,7 +393,7 @@ public:
   //! @rst
   //! Computes a device-wide segmented minimum using the less-than (``<``) operator.
   //!
-  //! - Uses ``std::numeric_limits<T>::max()`` as the initial value of the reduction for each segment.
+  //! - Uses ``::cuda::std::numeric_limits<T>::max()`` as the initial value of the reduction for each segment.
   //! - When input a contiguous sequence of segments, a single sequence
   //!   ``segment_offsets`` (of length ``num_segments + 1``) can be aliased for both
   //!   the ``d_begin_offsets`` and ``d_end_offsets`` parameters (where the latter is
@@ -508,8 +509,7 @@ public:
       d_begin_offsets,
       d_end_offsets,
       ::cuda::minimum<>{},
-      // TODO(bgruber): replace with ::cuda::std::numeric_limits<T>::max() (breaking change)
-      Traits<InputT>::Max(),
+      ::cuda::std::numeric_limits<InputT>::max(),
       stream);
   }
 
@@ -522,7 +522,7 @@ public:
   //!
   //!   - The minimum of the *i*\ :sup:`th` segment is written to
   //!     ``d_out[i].value`` and its offset in that segment is written to ``d_out[i].key``.
-  //!   - The ``{1, std::numeric_limits<T>::max()}`` tuple is produced for zero-length inputs
+  //!   - The ``{1, ::cuda::std::numeric_limits<T>::max()}`` tuple is produced for zero-length inputs
   //!
   //! - When input a contiguous sequence of segments, a single sequence
   //!   ``segment_offsets`` (of length ``num_segments + 1``) can be aliased for both
@@ -636,8 +636,7 @@ public:
     ArgIndexInputIteratorT d_indexed_in(d_in);
 
     // Initial value
-    // TODO Address https://github.com/NVIDIA/cub/issues/651
-    InitT initial_value{AccumT(1, Traits<InputValueT>::Max())};
+    InitT initial_value{AccumT(1, ::cuda::std::numeric_limits<InputValueT>::max())};
 
     using integral_offset_check = ::cuda::std::is_integral<OffsetT>;
     static_assert(integral_offset_check::value, "Offset iterator value type should be integral.");
@@ -666,7 +665,7 @@ public:
   //! @rst
   //! Computes a device-wide segmented maximum using the greater-than (``>``) operator.
   //!
-  //! - Uses ``std::numeric_limits<T>::lowest()`` as the initial value of the reduction.
+  //! - Uses ``::cuda::std::numeric_limits<T>::lowest()`` as the initial value of the reduction.
   //! - When input a contiguous sequence of segments, a single sequence
   //!   ``segment_offsets`` (of length ``num_segments + 1``) can be aliased
   //!   for both the ``d_begin_offsets`` and ``d_end_offsets`` parameters (where
@@ -771,8 +770,7 @@ public:
       d_begin_offsets,
       d_end_offsets,
       ::cuda::maximum<>{},
-      // TODO(bgruber): replace with ::cuda::std::numeric_limits<T>::lowest() (breaking change)
-      Traits<InputT>::Lowest(),
+      ::cuda::std::numeric_limits<InputT>::lowest(),
       stream);
   }
 
@@ -785,7 +783,7 @@ public:
   //!
   //!   - The maximum of the *i*\ :sup:`th` segment is written to
   //!     ``d_out[i].value`` and its offset in that segment is written to ``d_out[i].key``.
-  //!   - The ``{1, std::numeric_limits<T>::lowest()}`` tuple is produced for zero-length inputs
+  //!   - The ``{1, ::cuda::std::numeric_limits<T>::lowest()}`` tuple is produced for zero-length inputs
   //!
   //! - When input a contiguous sequence of segments, a single sequence
   //!   ``segment_offsets`` (of length ``num_segments + 1``) can be aliased
@@ -902,8 +900,7 @@ public:
     ArgIndexInputIteratorT d_indexed_in(d_in);
 
     // Initial value
-    // TODO Address https://github.com/NVIDIA/cub/issues/651
-    InitT initial_value{AccumT(1, Traits<InputValueT>::Lowest())};
+    InitT initial_value{AccumT(1, ::cuda::std::numeric_limits<InputValueT>::lowest())};
 
     using integral_offset_check = ::cuda::std::is_integral<OffsetT>;
     static_assert(integral_offset_check::value, "Offset iterator value type should be integral.");
diff --git a/cub/cub/util_type.cuh b/cub/cub/util_type.cuh
index 916cf6571ba..04d0062f048 100644
--- a/cub/cub/util_type.cuh
+++ b/cub/cub/util_type.cuh
@@ -780,18 +780,21 @@ enum Category
   FLOATING_POINT
 };
 
-/**
- * \brief Basic type traits
- */
-template <Category _CATEGORY, typename _UnsignedBits, typename T>
+namespace detail
+{
+struct is_primitive_impl;
+
+template <Category _CATEGORY, bool _PRIMITIVE, typename _UnsignedBits, typename T>
 struct BaseTraits
-{};
+{
+private:
+  friend struct is_primitive_impl;
+
+  static constexpr bool is_primitive = _PRIMITIVE;
+};
 
-/**
- * Basic type traits (unsigned primitive specialization)
- */
 template <typename _UnsignedBits, typename T>
-struct BaseTraits<UNSIGNED_INTEGER, _UnsignedBits, T>
+struct BaseTraits<UNSIGNED_INTEGER, true, _UnsignedBits, T>
 {
   static_assert(::cuda::std::numeric_limits<T>::is_specialized,
                 "Please also specialize cuda::std::numeric_limits for T");
@@ -810,6 +813,8 @@ struct BaseTraits<UNSIGNED_INTEGER, _UnsignedBits, T>
     return key;
   }
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::max()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Max()
   {
     UnsignedBits retval_bits = MAX_KEY;
@@ -818,6 +823,8 @@ struct BaseTraits<UNSIGNED_INTEGER, _UnsignedBits, T>
     return retval;
   }
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::lowest()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Lowest()
   {
     UnsignedBits retval_bits = LOWEST_KEY;
@@ -825,13 +832,15 @@ struct BaseTraits<UNSIGNED_INTEGER, _UnsignedBits, T>
     memcpy(&retval, &retval_bits, sizeof(T));
     return retval;
   }
+
+private:
+  friend struct is_primitive_impl;
+
+  static constexpr bool is_primitive = true;
 };
 
-/**
- * Basic type traits (signed primitive specialization)
- */
 template <typename _UnsignedBits, typename T>
-struct BaseTraits<SIGNED_INTEGER, _UnsignedBits, T>
+struct BaseTraits<SIGNED_INTEGER, true, _UnsignedBits, T>
 {
   static_assert(::cuda::std::numeric_limits<T>::is_specialized,
                 "Please also specialize cuda::std::numeric_limits for T");
@@ -852,27 +861,35 @@ struct BaseTraits<SIGNED_INTEGER, _UnsignedBits, T>
     return key ^ HIGH_BIT;
   };
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::max()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Max()
   {
     UnsignedBits retval = MAX_KEY;
     return reinterpret_cast<T&>(retval);
   }
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::lowest()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Lowest()
   {
     UnsignedBits retval = LOWEST_KEY;
     return reinterpret_cast<T&>(retval);
   }
+
+private:
+  friend struct is_primitive_impl;
+
+  static constexpr bool is_primitive = true;
 };
 
-/**
- * Basic type traits (fp primitive specialization)
- */
 template <typename _UnsignedBits, typename T>
-struct BaseTraits<FLOATING_POINT, _UnsignedBits, T>
+struct BaseTraits<FLOATING_POINT, true, _UnsignedBits, T>
 {
   static_assert(::cuda::std::numeric_limits<T>::is_specialized,
                 "Please also specialize cuda::std::numeric_limits for T");
+  static_assert(::cuda::is_floating_point<T>::value, "Please also specialize cuda::is_floating_point for T");
+  static_assert(::cuda::is_floating_point_v<T>, "Please also specialize cuda::is_floating_point_v for T");
 
   using UnsignedBits = _UnsignedBits;
 
@@ -892,48 +909,66 @@ struct BaseTraits<FLOATING_POINT, _UnsignedBits, T>
     return key ^ mask;
   };
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::max()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Max()
   {
     return ::cuda::std::numeric_limits<T>::max();
   }
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::lowest()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Lowest()
   {
     return ::cuda::std::numeric_limits<T>::lowest();
   }
-};
 
-/**
- * \brief Numeric type traits
- */
-// clang-format off
-template <typename T> struct NumericTraits :            BaseTraits<NOT_A_NUMBER, T, T> {};
-
-template <> struct NumericTraits<NullType> :            BaseTraits<NOT_A_NUMBER, NullType, NullType> {};
-
-template <> struct NumericTraits<char> :                BaseTraits<(::cuda::std::numeric_limits<char>::is_signed) ? SIGNED_INTEGER : UNSIGNED_INTEGER, unsigned char, char> {};
-template <> struct NumericTraits<signed char> :         BaseTraits<SIGNED_INTEGER, unsigned char, signed char> {};
-template <> struct NumericTraits<short> :               BaseTraits<SIGNED_INTEGER, unsigned short, short> {};
-template <> struct NumericTraits<int> :                 BaseTraits<SIGNED_INTEGER, unsigned int, int> {};
-template <> struct NumericTraits<long> :                BaseTraits<SIGNED_INTEGER, unsigned long, long> {};
-template <> struct NumericTraits<long long> :           BaseTraits<SIGNED_INTEGER, unsigned long long, long long> {};
+private:
+  friend struct is_primitive_impl;
 
-template <> struct NumericTraits<unsigned char> :       BaseTraits<UNSIGNED_INTEGER, unsigned char, unsigned char> {};
-template <> struct NumericTraits<unsigned short> :      BaseTraits<UNSIGNED_INTEGER, unsigned short, unsigned short> {};
-template <> struct NumericTraits<unsigned int> :        BaseTraits<UNSIGNED_INTEGER, unsigned int, unsigned int> {};
-template <> struct NumericTraits<unsigned long> :       BaseTraits<UNSIGNED_INTEGER, unsigned long, unsigned long> {};
-template <> struct NumericTraits<unsigned long long> :  BaseTraits<UNSIGNED_INTEGER, unsigned long long, unsigned long long> {};
+  static constexpr bool is_primitive = true;
+};
+} // namespace detail
 
+//! Use this class as base when specializing \ref NumericTraits for primitive signed/unsigned integers or floating-point
+//! types.
+template <Category _CATEGORY, bool _PRIMITIVE, typename _UnsignedBits, typename T>
+using BaseTraits = detail::BaseTraits<_CATEGORY, _PRIMITIVE, _UnsignedBits, T>;
+
+//! Numeric type traits for radix sort key operations, decoupled lookback and tuning. You can specialize this template
+//! for your own types if:
+//! * There is an unsigned integral type of equal size
+//! * The size of the type is smaller than 64bits
+//! * The arithmetic throughput of the type is similar to other built-in types of the same size
+//! For other types, if you want to use them with radix sort, please use the decomposer interface of the radix sort.
+// clang-format off
+template <typename T> struct NumericTraits :            BaseTraits<NOT_A_NUMBER, false, T, T> {};
+
+template <> struct NumericTraits<NullType> :            BaseTraits<NOT_A_NUMBER, false, NullType, NullType> {};
+
+template <> struct NumericTraits<char> :                BaseTraits<(::cuda::std::numeric_limits<char>::is_signed) ? SIGNED_INTEGER : UNSIGNED_INTEGER, true, unsigned char, char> {};
+template <> struct NumericTraits<signed char> :         BaseTraits<SIGNED_INTEGER, true, unsigned char, signed char> {};
+template <> struct NumericTraits<short> :               BaseTraits<SIGNED_INTEGER, true, unsigned short, short> {};
+template <> struct NumericTraits<int> :                 BaseTraits<SIGNED_INTEGER, true, unsigned int, int> {};
+template <> struct NumericTraits<long> :                BaseTraits<SIGNED_INTEGER, true, unsigned long, long> {};
+template <> struct NumericTraits<long long> :           BaseTraits<SIGNED_INTEGER, true, unsigned long long, long long> {};
+
+template <> struct NumericTraits<unsigned char> :       BaseTraits<UNSIGNED_INTEGER, true, unsigned char, unsigned char> {};
+template <> struct NumericTraits<unsigned short> :      BaseTraits<UNSIGNED_INTEGER, true, unsigned short, unsigned short> {};
+template <> struct NumericTraits<unsigned int> :        BaseTraits<UNSIGNED_INTEGER, true, unsigned int, unsigned int> {};
+template <> struct NumericTraits<unsigned long> :       BaseTraits<UNSIGNED_INTEGER, true, unsigned long, unsigned long> {};
+template <> struct NumericTraits<unsigned long long> :  BaseTraits<UNSIGNED_INTEGER, true, unsigned long long, unsigned long long> {};
+// clang-format on
 
-#if _CCCL_HAS_INT128()
+#  if _CCCL_HAS_INT128()
 template <>
 struct NumericTraits<__uint128_t>
 {
-  using T = __uint128_t;
+  using T            = __uint128_t;
   using UnsignedBits = __uint128_t;
 
-  static constexpr UnsignedBits   LOWEST_KEY  = UnsignedBits(0);
-  static constexpr UnsignedBits   MAX_KEY     = UnsignedBits(-1);
+  static constexpr UnsignedBits LOWEST_KEY = UnsignedBits(0);
+  static constexpr UnsignedBits MAX_KEY    = UnsignedBits(-1);
 
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits TwiddleIn(UnsignedBits key)
   {
@@ -945,26 +980,35 @@ struct NumericTraits<__uint128_t>
     return key;
   }
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::max()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Max()
   {
     return MAX_KEY;
   }
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::lowest()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Lowest()
   {
     return LOWEST_KEY;
   }
+
+private:
+  friend struct detail::is_primitive_impl;
+
+  static constexpr bool is_primitive = false;
 };
 
 template <>
 struct NumericTraits<__int128_t>
 {
-  using T = __int128_t;
+  using T            = __int128_t;
   using UnsignedBits = __uint128_t;
 
-  static constexpr UnsignedBits   HIGH_BIT    = UnsignedBits(1) << ((sizeof(UnsignedBits) * 8) - 1);
-  static constexpr UnsignedBits   LOWEST_KEY  = HIGH_BIT;
-  static constexpr UnsignedBits   MAX_KEY     = UnsignedBits(-1) ^ HIGH_BIT;
+  static constexpr UnsignedBits HIGH_BIT   = UnsignedBits(1) << ((sizeof(UnsignedBits) * 8) - 1);
+  static constexpr UnsignedBits LOWEST_KEY = HIGH_BIT;
+  static constexpr UnsignedBits MAX_KEY    = UnsignedBits(-1) ^ HIGH_BIT;
 
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits TwiddleIn(UnsignedBits key)
   {
@@ -976,83 +1020,82 @@ struct NumericTraits<__int128_t>
     return key ^ HIGH_BIT;
   };
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::max()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Max()
   {
     UnsignedBits retval = MAX_KEY;
     return reinterpret_cast<T&>(retval);
   }
 
+  //! deprecated [Since 3.0]
+  CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits<T>::lowest()")
   static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Lowest()
   {
     UnsignedBits retval = LOWEST_KEY;
     return reinterpret_cast<T&>(retval);
   }
+
+private:
+  friend struct detail::is_primitive_impl;
+
+  static constexpr bool is_primitive = false;
 };
-#endif // _CCCL_HAS_INT128()
+#  endif // _CCCL_HAS_INT128()
 
-template <> struct NumericTraits<float> :               BaseTraits<FLOATING_POINT,unsigned int, float> {};
-template <> struct NumericTraits<double> :              BaseTraits<FLOATING_POINT,unsigned long long, double> {};
+// clang-format off
+template <> struct NumericTraits<float> :               BaseTraits<FLOATING_POINT, true, unsigned int, float> {};
+template <> struct NumericTraits<double> :              BaseTraits<FLOATING_POINT, true, unsigned long long, double> {};
 #  if _CCCL_HAS_NVFP16()
-    template <> struct NumericTraits<__half> :          BaseTraits<FLOATING_POINT,unsigned short, __half> {};
+    template <> struct NumericTraits<__half> :          BaseTraits<FLOATING_POINT, true, unsigned short, __half> {};
 #  endif // _CCCL_HAS_NVFP16()
 #  if _CCCL_HAS_NVBF16()
-    template <> struct NumericTraits<__nv_bfloat16> :   BaseTraits<FLOATING_POINT,unsigned short, __nv_bfloat16> {};
+    template <> struct NumericTraits<__nv_bfloat16> :   BaseTraits<FLOATING_POINT, true, unsigned short, __nv_bfloat16> {};
 #  endif // _CCCL_HAS_NVBF16()
 
 #if _CCCL_HAS_NVFP8()
-    template <> struct NumericTraits<__nv_fp8_e4m3> :   BaseTraits<FLOATING_POINT, __nv_fp8_storage_t, __nv_fp8_e4m3> {};
-    template <> struct NumericTraits<__nv_fp8_e5m2> :   BaseTraits<FLOATING_POINT, __nv_fp8_storage_t, __nv_fp8_e5m2> {};
+    template <> struct NumericTraits<__nv_fp8_e4m3> :   BaseTraits<FLOATING_POINT, true, __nv_fp8_storage_t, __nv_fp8_e4m3> {};
+    template <> struct NumericTraits<__nv_fp8_e5m2> :   BaseTraits<FLOATING_POINT, true, __nv_fp8_storage_t, __nv_fp8_e5m2> {};
 #endif // _CCCL_HAS_NVFP8()
 
-template <> struct NumericTraits<bool> :                BaseTraits<UNSIGNED_INTEGER, typename UnitWord<bool>::VolatileWord, bool> {};
+template <> struct NumericTraits<bool> :                BaseTraits<UNSIGNED_INTEGER, true, typename UnitWord<bool>::VolatileWord, bool> {};
 // clang-format on
 
-/**
- * \brief Type traits
- */
+namespace detail
+{
 template <typename T>
 struct Traits : NumericTraits<::cuda::std::remove_cv_t<T>>
 {};
+} // namespace detail
+
+//! \brief Query type traits for radix sort key operations, decoupled lookback and tunings. To add support for your own
+//! primitive types please specialize \ref NumericTraits.
+template <typename T>
+using Traits = detail::Traits<T>;
 
 namespace detail
 {
-// __uint128_t and __int128_t are not primitive
+// we cannot befriend is_primitive on GCC < 11, since it's a template (bug)
+struct is_primitive_impl
+{
+  // must be a struct instead of an alias, so the access of Traits<T>::is_primitive happens in the context of this class
+  template <typename T>
+  struct is_primitive : ::cuda::std::bool_constant<Traits<T>::is_primitive>
+  {};
+};
+// This trait serves two purposes:
+// 1. It is used for tunings to detect whether we have a build-in arithmetic type for which we can expect certain
+// arithmetic throughput. E.g.: we expect all primitive types of the same size to show roughly similar performance.
+// 2. Decoupled lookback uses this trait to determine whether there is a machine word twice the size of T which can be
+// loaded/stored with a single instruction.
+// TODO(bgruber): for 2. we should probably just check whether sizeof(T) * 2 <= sizeof(int128) (or 256-bit on SM100)
+// Users must be able to hook into both scenarios with their custom types, so this trait must depend on cub::Traits
 template <typename T>
-using is_primitive = ::cuda::std::bool_constant<is_one_of<
-  T,
-  char,
-  signed char,
-  short,
-  int,
-  long,
-  long long,
-  unsigned char,
-  unsigned short,
-  unsigned int,
-  unsigned long,
-  unsigned long long,
-  bool,
-  float,
-  double
-#  if _CCCL_HAS_NVFP16()
-  ,
-  __half
-#  endif // _CCCL_HAS_NVFP16()
-#  if _CCCL_HAS_NVBF16()
-  ,
-  __nv_bfloat16
-#  endif // _CCCL_HAS_NVBF16()
-#  if _CCCL_HAS_NVFP8()
-  ,
-  __nv_fp8_e4m3,
-  __nv_fp8_e5m2
-#  endif // _CCCL_HAS_NVFP8()
-  >()>;
-
-#  ifndef _CCCL_NO_VARIABLE_TEMPLATES
+struct is_primitive : is_primitive_impl::is_primitive<T>
+{};
+
 template <typename T>
-inline constexpr bool is_primitive_v = is_primitive<T>::value;
-#  endif // !_CCCL_NO_VARIABLE_TEMPLATES
+_CCCL_INLINE_VAR constexpr bool is_primitive_v = is_primitive<T>::value;
 } // namespace detail
 
 #endif // _CCCL_DOXYGEN_INVOKED
diff --git a/cub/examples/device/example_device_radix_sort_custom.cu b/cub/examples/device/example_device_radix_sort_custom.cu
index d0bcf920fbb..00a892f40f0 100644
--- a/cub/examples/device/example_device_radix_sort_custom.cu
+++ b/cub/examples/device/example_device_radix_sort_custom.cu
@@ -114,10 +114,10 @@ int main()
   std::cout << "l:\t" << to_binary_representation(l) << '\n';
   std::cout << "g:\t" << to_binary_representation(g) << "\n\n";
 
-  std::cout << "As you can see, `l` key happened to be larger in the bit-lexicographicl order.\n";
-  std::cout << "Since there's no reflection in C++, we can't inspect the type and convert \n";
-  std::cout << "each field into the bit-lexicographicl order. You can tell CUB how to do that\n";
-  std::cout << "by specializing cub::RadixTraits for the `custom_t`:\n\n";
+  std::cout << "As you can see, `l` key happened to be larger in the bit-lexicographical order.\n";
+  std::cout << "Since there's no reflection in C++ (yet), we can't inspect the type and convert \n";
+  std::cout << "each field into the bit-lexicographical order. You can tell CUB how to do that\n";
+  std::cout << "by providing a decomposer for the `custom_t`:\n\n";
 
   std::cout << "\tstruct decomposer_t \n";
   std::cout << "\t{\n";
@@ -132,7 +132,7 @@ int main()
   std::cout << "Decomposer allows you to specify which fields are most significant and which\n";
   std::cout << "are least significant. In our case, `f` is the most significant field and\n";
   std::cout << "`i` is the least significant field. The decomposer is then used by CUB to convert\n";
-  std::cout << "the `custom_t` into the bit-lexicographicl order:\n\n";
+  std::cout << "the `custom_t` into the bit-lexicographical order:\n\n";
 
   using conversion_policy = cub::detail::radix::traits_t<custom_t>::bit_ordered_conversion_policy;
   l                       = conversion_policy::to_bit_ordered(decomposer_t{}, l);
@@ -148,7 +148,7 @@ int main()
   std::cout << "g:\t" << to_binary_representation(g) << "\n\n";
 
   std::cout << '\n';
-  std::cout << "As you can see, `g` is now actually larger than `l` in the bit-lexicographicl order.\n";
+  std::cout << "As you can see, `g` is now actually larger than `l` in the bit-lexicographical order.\n";
   std::cout << "After binning, CUB is able to restore the original key:\n\n";
 
   l = conversion_policy::from_bit_ordered(decomposer_t{}, l);
diff --git a/cub/test/catch2_test_device_reduce.cuh b/cub/test/catch2_test_device_reduce.cuh
index 5760aa93536..a33a2a02439 100644
--- a/cub/test/catch2_test_device_reduce.cuh
+++ b/cub/test/catch2_test_device_reduce.cuh
@@ -114,29 +114,6 @@ CUB_NAMESPACE_END
 
 CUB_NAMESPACE_BEGIN
 
-// TODO(bgruber): drop this when we drop cub::Traits
-template <template <typename> class... Policies>
-struct NumericTraits<c2h::custom_type_t<Policies...>>
-{
-  using custom_t = c2h::custom_type_t<Policies...>;
-
-  __host__ __device__ static custom_t Max()
-  {
-    custom_t val{};
-    val.key = NumericTraits<decltype(std::declval<custom_t>().key)>::Max();
-    val.val = NumericTraits<decltype(std::declval<custom_t>().val)>::Max();
-    return val;
-  }
-
-  __host__ __device__ static custom_t Lowest()
-  {
-    custom_t val{};
-    val.key = NumericTraits<decltype(std::declval<custom_t>().key)>::Lowest();
-    val.val = NumericTraits<decltype(std::declval<custom_t>().val)>::Lowest();
-    return val;
-  }
-};
-
 template <typename Key, typename Value>
 static std::ostream& operator<<(std::ostream& os, const KeyValuePair<Key, Value>& val)
 {
diff --git a/cub/test/catch2_test_device_segmented_sort_keys.cu b/cub/test/catch2_test_device_segmented_sort_keys.cu
index eb6cbd5eeca..01a28a5699b 100644
--- a/cub/test/catch2_test_device_segmented_sort_keys.cu
+++ b/cub/test/catch2_test_device_segmented_sort_keys.cu
@@ -32,7 +32,9 @@
 
 #include "catch2_radix_sort_helper.cuh"
 #include "catch2_segmented_sort_helper.cuh"
+#include <c2h/bfloat16.cuh>
 #include <c2h/catch2_test_helper.h>
+#include <c2h/half.cuh>
 
 // FIXME: Graph launch disabled, algorithm syncs internally. WAR exists for device-launch, figure out how to enable for
 // graph launch.
diff --git a/cub/test/catch2_test_util_type.cu b/cub/test/catch2_test_util_type.cu
index b3514fa2e79..bb34605935b 100644
--- a/cub/test/catch2_test_util_type.cu
+++ b/cub/test/catch2_test_util_type.cu
@@ -80,3 +80,30 @@ C2H_TEST("Test CUB_DEFINE_DETECT_NESTED_TYPE", "[util][type]")
   STATIC_REQUIRE(cat_detect<HasCat>::value);
   STATIC_REQUIRE(!cat_detect<HasDog>::value);
 }
+
+// Lots of libraries (like pytorch or tensorflow) bring their own half types and customize cub::Traits for them.
+struct CustomHalf
+{
+  int16_t payload;
+};
+
+C2H_TEST("Test CustomHalf", "[util][type]")
+{
+  // type not registered with cub::Traits
+  STATIC_REQUIRE(!cub::detail::is_primitive<CustomHalf>::value);
+  STATIC_REQUIRE(!cuda::std::is_floating_point<CustomHalf>::value);
+  STATIC_REQUIRE(!cuda::std::is_floating_point_v<CustomHalf>);
+  STATIC_REQUIRE(!cuda::is_floating_point<CustomHalf>::value);
+  STATIC_REQUIRE(!cuda::is_floating_point_v<CustomHalf>);
+  STATIC_REQUIRE(!cuda::std::numeric_limits<CustomHalf>::is_specialized);
+
+  // type registered with cub::Traits (specializes NumericTraits, numeric_limits, and is_floating_point)
+  STATIC_REQUIRE(cub::detail::is_primitive<half_t>::value);
+  STATIC_REQUIRE(!cuda::std::is_floating_point<half_t>::value); // the std traits are not affected
+  STATIC_REQUIRE(!cuda::std::is_floating_point_v<half_t>);
+  STATIC_REQUIRE(cuda::is_floating_point<half_t>::value);
+  STATIC_REQUIRE(cuda::is_floating_point_v<half_t>);
+  STATIC_REQUIRE(cuda::std::numeric_limits<half_t>::is_specialized);
+  CHECK(cuda::std::numeric_limits<half_t>::max() == half_t::max());
+  CHECK(cuda::std::numeric_limits<half_t>::lowest() == half_t::lowest());
+}
diff --git a/cub/test/test_util.h b/cub/test/test_util.h
index 4041af6841e..88552be32ac 100644
--- a/cub/test/test_util.h
+++ b/cub/test/test_util.h
@@ -966,31 +966,6 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, TestFoo& va
   InitValue(gen_mode, value.w, index);
 }
 
-/// numeric_limits<TestFoo> specialization
-CUB_NAMESPACE_BEGIN
-template <>
-struct NumericTraits<TestFoo>
-{
-  __host__ __device__ static TestFoo Max()
-  {
-    return TestFoo::MakeTestFoo(
-      NumericTraits<long long>::Max(),
-      NumericTraits<int>::Max(),
-      NumericTraits<short>::Max(),
-      NumericTraits<char>::Max());
-  }
-
-  __host__ __device__ static TestFoo Lowest()
-  {
-    return TestFoo::MakeTestFoo(
-      NumericTraits<long long>::Lowest(),
-      NumericTraits<int>::Lowest(),
-      NumericTraits<short>::Lowest(),
-      NumericTraits<char>::Lowest());
-  }
-};
-CUB_NAMESPACE_END
-
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 template <>
 class numeric_limits<TestFoo>
@@ -1121,23 +1096,6 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, TestBar& va
   InitValue(gen_mode, value.y, index);
 }
 
-/// numeric_limits<TestBar> specialization
-CUB_NAMESPACE_BEGIN
-template <>
-struct NumericTraits<TestBar>
-{
-  __host__ __device__ static TestBar Max()
-  {
-    return TestBar(NumericTraits<long long>::Max(), NumericTraits<int>::Max());
-  }
-
-  __host__ __device__ static TestBar Lowest()
-  {
-    return TestBar(NumericTraits<long long>::Lowest(), NumericTraits<int>::Lowest());
-  }
-};
-CUB_NAMESPACE_END
-
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 template <>
 class numeric_limits<TestBar>

From 6c15254cf281a1faacd741a0e957bacff2bc60d9 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Mon, 24 Feb 2025 16:54:44 +0100
Subject: [PATCH 02/17] Define is_floating_point_v in terms of
 is_floating_point (#3923)

---
 .../include/cuda/__type_traits/is_floating_point.h  | 13 +++++--------
 1 file changed, 5 insertions(+), 8 deletions(-)

diff --git a/libcudacxx/include/cuda/__type_traits/is_floating_point.h b/libcudacxx/include/cuda/__type_traits/is_floating_point.h
index e253315a672..8d7fe63bc6f 100644
--- a/libcudacxx/include/cuda/__type_traits/is_floating_point.h
+++ b/libcudacxx/include/cuda/__type_traits/is_floating_point.h
@@ -27,18 +27,15 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_CUDA
 
-template <class _Tp>
-struct _CCCL_TYPE_VISIBILITY_DEFAULT is_floating_point
-    : _CUDA_VSTD::bool_constant<_CUDA_VSTD::is_floating_point<_CUDA_VSTD::remove_cv_t<_Tp>>::value
-                                || _CUDA_VSTD::__is_extended_floating_point<_CUDA_VSTD::remove_cv_t<_Tp>>::value>
-{};
-
-#if !defined(_CCCL_NO_VARIABLE_TEMPLATES)
+//! Tells whether a type is a floating point type. Users are allowed this template for their own types.
 template <class _Tp>
 _CCCL_INLINE_VAR constexpr bool is_floating_point_v =
   _CUDA_VSTD::is_floating_point_v<_CUDA_VSTD::remove_cv_t<_Tp>>
   || _CUDA_VSTD::__is_extended_floating_point_v<_CUDA_VSTD::remove_cv_t<_Tp>>;
-#endif // !_CCCL_NO_VARIABLE_TEMPLATES
+
+// we define the trait as alias, so users cannot specialize it (they should specialize the variable template instead)
+template <class _Tp>
+using is_floating_point = _CUDA_VSTD::bool_constant<is_floating_point_v<_Tp>>;
 
 _LIBCUDACXX_END_NAMESPACE_CUDA
 

From c065aa84b285a28c8e753f87081bbc4c6e6edf1f Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Mon, 24 Feb 2025 18:10:38 +0100
Subject: [PATCH 03/17] Replace usage of CUB_MIN|MAX in reduce (#3927)

---
 cub/cub/agent/agent_reduce.cuh                         |  2 +-
 .../block_reduce_raking_commutative_only.cuh           |  2 +-
 .../specializations/block_reduce_warp_reductions.cuh   |  2 +-
 cub/cub/device/dispatch/dispatch_reduce_by_key.cuh     |  4 ++--
 .../device/dispatch/tuning/tuning_reduce_by_key.cuh    |  8 ++++----
 cub/cub/warp/specializations/warp_reduce_smem.cuh      |  2 +-
 cub/test/catch2_test_device_reduce.cuh                 | 10 ++++++----
 7 files changed, 16 insertions(+), 14 deletions(-)

diff --git a/cub/cub/agent/agent_reduce.cuh b/cub/cub/agent/agent_reduce.cuh
index e6610e52c14..c6f7c0a95a4 100644
--- a/cub/cub/agent/agent_reduce.cuh
+++ b/cub/cub/agent/agent_reduce.cuh
@@ -159,7 +159,7 @@ struct AgentReduce
   static constexpr int BLOCK_THREADS      = AgentReducePolicy::BLOCK_THREADS;
   static constexpr int ITEMS_PER_THREAD   = AgentReducePolicy::ITEMS_PER_THREAD;
   static constexpr int TILE_ITEMS         = BLOCK_THREADS * ITEMS_PER_THREAD;
-  static constexpr int VECTOR_LOAD_LENGTH = CUB_MIN(ITEMS_PER_THREAD, AgentReducePolicy::VECTOR_LOAD_LENGTH);
+  static constexpr int VECTOR_LOAD_LENGTH = _CUDA_VSTD::min(ITEMS_PER_THREAD, AgentReducePolicy::VECTOR_LOAD_LENGTH);
 
   // Can vectorize according to the policy if the input iterator is a native
   // pointer to a primitive type
diff --git a/cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh b/cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh
index 28ff55b5fe0..6f962499984 100644
--- a/cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh
+++ b/cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh
@@ -96,7 +96,7 @@ struct BlockReduceRakingCommutativeOnly
     RAKING_THREADS = WARP_THREADS,
 
     /// Number of threads actually sharing items with the raking threads
-    SHARING_THREADS = CUB_MAX(1, BLOCK_THREADS - RAKING_THREADS),
+    SHARING_THREADS = _CUDA_VSTD::max(1, BLOCK_THREADS - RAKING_THREADS),
 
     /// Number of raking elements per warp synchronous raking thread
     SEGMENT_LENGTH = SHARING_THREADS / WARP_THREADS,
diff --git a/cub/cub/block/specializations/block_reduce_warp_reductions.cuh b/cub/cub/block/specializations/block_reduce_warp_reductions.cuh
index 76de43da1d7..07faa04dca0 100644
--- a/cub/cub/block/specializations/block_reduce_warp_reductions.cuh
+++ b/cub/cub/block/specializations/block_reduce_warp_reductions.cuh
@@ -84,7 +84,7 @@ struct BlockReduceWarpReductions
     WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS,
 
     /// The logical warp size for warp reductions
-    LOGICAL_WARP_SIZE = CUB_MIN(BLOCK_THREADS, WARP_THREADS),
+    LOGICAL_WARP_SIZE = _CUDA_VSTD::min(BLOCK_THREADS, WARP_THREADS),
 
     /// Whether or not the logical warp size evenly divides the thread block size
     EVEN_WARP_MULTIPLE = (BLOCK_THREADS % LOGICAL_WARP_SIZE == 0)
diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
index f3eb01d1e25..439dbd48634 100644
--- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
+++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
@@ -344,7 +344,7 @@ struct DispatchReduceByKey
       }
 
       // Log init_kernel configuration
-      int init_grid_size = CUB_MAX(1, ::cuda::ceil_div(num_tiles, INIT_KERNEL_THREADS));
+      int init_grid_size = _CUDA_VSTD::max(1, ::cuda::ceil_div(num_tiles, INIT_KERNEL_THREADS));
 
 #ifdef CUB_DEBUG_LOG
       _CubLog("Invoking init_kernel<<<%d, %d, 0, %lld>>>()\n", init_grid_size, INIT_KERNEL_THREADS, (long long) stream);
@@ -392,7 +392,7 @@ struct DispatchReduceByKey
       }
 
       // Run grids in epochs (in case number of tiles exceeds max x-dimension
-      int scan_grid_size = CUB_MIN(num_tiles, max_dim_x);
+      int scan_grid_size = _CUDA_VSTD::min(num_tiles, max_dim_x);
       for (int start_tile = 0; start_tile < num_tiles; start_tile += scan_grid_size)
       {
 // Log reduce_by_key_kernel configuration
diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
index 7d7e4ecd738..321749532e2 100644
--- a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
@@ -45,7 +45,8 @@
 #include <cub/util_device.cuh>
 #include <cub/util_type.cuh>
 
-#include <cuda/std/__algorithm/max.h>
+#include <cuda/cmath>
+#include <cuda/std/__algorithm/clamp.h>
 
 CUB_NAMESPACE_BEGIN
 
@@ -871,9 +872,8 @@ struct policy_hub
     static constexpr int items_per_thread =
       (max_input_bytes <= 8)
         ? 6
-        // TODO(bgruber): use ceil_div and clamp in C++14
-        : CUB_MIN(nominal_4B_items_per_thread,
-                  CUB_MAX(1, ((nominal_4B_items_per_thread * 8) + combined_input_bytes - 1) / combined_input_bytes));
+        : ::cuda::std::clamp(
+            ::cuda::ceil_div(nominal_4B_items_per_thread * 8, combined_input_bytes), 1, nominal_4B_items_per_thread);
 
     using ReduceByKeyPolicyT =
       AgentReduceByKeyPolicy<128,
diff --git a/cub/cub/warp/specializations/warp_reduce_smem.cuh b/cub/cub/warp/specializations/warp_reduce_smem.cuh
index 02b8295460a..7343909d924 100644
--- a/cub/cub/warp/specializations/warp_reduce_smem.cuh
+++ b/cub/cub/warp/specializations/warp_reduce_smem.cuh
@@ -241,7 +241,7 @@ struct WarpReduceSmem
     // Clip the next segment at the warp boundary if necessary
     if (LOGICAL_WARP_THREADS != 32)
     {
-      next_flag = CUB_MIN(next_flag, LOGICAL_WARP_THREADS);
+      next_flag = _CUDA_VSTD::min(next_flag, LOGICAL_WARP_THREADS);
     }
 
 #pragma unroll
diff --git a/cub/test/catch2_test_device_reduce.cuh b/cub/test/catch2_test_device_reduce.cuh
index a33a2a02439..c43468a78b7 100644
--- a/cub/test/catch2_test_device_reduce.cuh
+++ b/cub/test/catch2_test_device_reduce.cuh
@@ -56,9 +56,10 @@ template <>
 _LIBCUDACXX_HIDE_FROM_ABI __half minimum<void>::operator()<__half, __half>(const __half& a, const __half& b) const
 {
 #  if defined(__CUDA_NO_HALF_OPERATORS__)
-  return CUB_MIN(__half2float(a), __half2float(b));
+  return ::cuda::std::min(__half2float(a), __half2float(b));
 #  else // ^^^ __CUDA_NO_HALF_OPERATORS__ ^^^ / vvv !__CUDA_NO_HALF_OPERATORS__ vvv
-  NV_IF_TARGET(NV_PROVIDES_SM_53, (return CUB_MIN(a, b);), (return CUB_MIN(__half2float(a), __half2float(b));));
+  NV_IF_TARGET(
+    NV_PROVIDES_SM_53, (return ::cuda::std::min(a, b);), (return ::cuda::std::min(__half2float(a), __half2float(b));));
 #  endif // !__CUDA_NO_HALF_OPERATORS__
 }
 
@@ -66,9 +67,10 @@ template <>
 _LIBCUDACXX_HIDE_FROM_ABI __half maximum<void>::operator()<__half, __half>(const __half& a, const __half& b) const
 {
 #  if defined(__CUDA_NO_HALF_OPERATORS__)
-  return CUB_MAX(__half2float(a), __half2float(b));
+  return ::cuda::std::max(__half2float(a), __half2float(b));
 #  else // ^^^ __CUDA_NO_HALF_OPERATORS__ ^^^ / vvv !__CUDA_NO_HALF_OPERATORS__ vvv
-  NV_IF_TARGET(NV_PROVIDES_SM_53, (return CUB_MAX(a, b);), (return CUB_MAX(__half2float(a), __half2float(b));));
+  NV_IF_TARGET(
+    NV_PROVIDES_SM_53, (return ::cuda::std::max(a, b);), (return ::cuda::std::max(__half2float(a), __half2float(b));));
 #  endif // !__CUDA_NO_HALF_OPERATORS__
 }
 

From 8486ea8031e47e45ffdd6b251ba786a33d6bce95 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Mon, 24 Feb 2025 20:41:33 +0100
Subject: [PATCH 04/17] Move to cuda::std::iterator_traits in CUB (#3924)

---
 c2h/include/c2h/fill_striped.h                |  2 +-
 cub/cub/agent/agent_batch_memcpy.cuh          | 12 ++--
 cub/cub/agent/agent_histogram.cuh             |  2 +-
 cub/cub/agent/agent_merge.cuh                 |  4 +-
 cub/cub/agent/agent_reduce.cuh                |  2 +-
 cub/cub/agent/agent_reduce_by_key.cuh         |  4 +-
 cub/cub/agent/agent_rle.cuh                   |  2 +-
 cub/cub/agent/agent_scan.cuh                  |  2 +-
 cub/cub/agent/agent_scan_by_key.cuh           |  4 +-
 cub/cub/agent/agent_select_if.cuh             |  4 +-
 cub/cub/agent/agent_three_way_partition.cuh   |  2 +-
 cub/cub/agent/agent_unique_by_key.cuh         |  4 +-
 cub/cub/block/block_load.cuh                  |  2 +-
 cub/cub/block/block_merge_sort.cuh            |  4 +-
 cub/cub/block/block_store.cuh                 |  2 +-
 cub/cub/device/device_for.cuh                 | 16 +++---
 cub/cub/device/device_histogram.cuh           | 12 ++--
 cub/cub/device/device_memcpy.cuh              |  4 +-
 cub/cub/device/device_reduce.cuh              | 14 ++---
 cub/cub/device/device_run_length_encode.cuh   |  2 +-
 cub/cub/device/device_scan.cuh                |  6 +-
 cub/cub/device/device_segmented_reduce.cuh    | 10 ++--
 .../dispatch/dispatch_adjacent_difference.cuh |  2 +-
 .../device/dispatch/dispatch_batch_memcpy.cuh | 19 ++++---
 .../device/dispatch/dispatch_histogram.cuh    | 10 ++--
 cub/cub/device/dispatch/dispatch_merge.cuh    |  4 +-
 .../device/dispatch/dispatch_merge_sort.cuh   | 12 ++--
 cub/cub/device/dispatch/dispatch_reduce.cuh   | 16 +++---
 .../dispatch/dispatch_reduce_by_key.cuh       |  8 +--
 cub/cub/device/dispatch/dispatch_rle.cuh      |  2 +-
 cub/cub/device/dispatch/dispatch_scan.cuh     |  8 +--
 .../device/dispatch/dispatch_scan_by_key.cuh  | 15 ++---
 .../device/dispatch/dispatch_select_if.cuh    |  4 +-
 .../dispatch/dispatch_three_way_partition.cuh |  2 +-
 .../dispatch/dispatch_unique_by_key.cuh       | 23 ++++----
 cub/cub/device/dispatch/kernels/transform.cuh |  4 +-
 .../tuning/tuning_adjacent_difference.cuh     |  2 +-
 .../dispatch/tuning/tuning_merge_sort.cuh     |  2 +-
 .../dispatch/tuning/tuning_scan_by_key.cuh    |  2 +-
 .../dispatch/tuning/tuning_transform.cuh      |  6 +-
 .../dispatch/tuning/tuning_unique_by_key.cuh  |  2 +-
 cub/cub/iterator/arg_index_input_iterator.cuh |  3 +-
 .../cache_modified_input_iterator.cuh         |  1 -
 .../cache_modified_output_iterator.cuh        |  1 -
 cub/cub/iterator/tex_obj_input_iterator.cuh   |  1 -
 cub/cub/thread/thread_load.cuh                |  6 +-
 cub/cub/thread/thread_search.cuh              |  2 +-
 cub/cub/util_type.cuh                         | 46 ++++++++-------
 cub/test/catch2_test_block_load.cu            |  2 +-
 cub/test/catch2_test_block_radix_sort.cuh     |  6 +-
 .../catch2_test_block_run_length_decode.cu    |  4 +-
 cub/test/catch2_test_block_store.cu           |  4 +-
 cub/test/catch2_test_device_reduce.cu         |  2 +-
 cub/test/catch2_test_device_reduce.cuh        |  2 +-
 cub/test/catch2_test_device_scan.cu           |  4 +-
 cub/test/catch2_test_device_scan.cuh          | 16 +++---
 cub/test/catch2_test_device_scan_by_key.cu    |  2 +-
 cub/test/catch2_test_device_scan_iterators.cu |  2 +-
 .../catch2_test_device_segmented_reduce.cu    |  2 +-
 cub/test/catch2_test_device_select_common.cuh |  4 +-
 cub/test/catch2_test_iterator.cu              | 56 ++++++++++---------
 cub/test/catch2_test_vsmem.cu                 |  6 +-
 thrust/thrust/system/cuda/detail/scan.h       |  2 +-
 63 files changed, 223 insertions(+), 210 deletions(-)

diff --git a/c2h/include/c2h/fill_striped.h b/c2h/include/c2h/fill_striped.h
index eab05478c88..a05b19f4bdf 100644
--- a/c2h/include/c2h/fill_striped.h
+++ b/c2h/include/c2h/fill_striped.h
@@ -139,7 +139,7 @@ struct scalar_to_vec_t
 template <int LogicalWarpThreads, int ItemsPerThread, int BlockThreads, typename IteratorT>
 void fill_striped(IteratorT it)
 {
-  using T = cub::detail::value_t<IteratorT>;
+  using T = cub::detail::it_value_t<IteratorT>;
 
   constexpr int warps_in_block = BlockThreads / LogicalWarpThreads;
   constexpr int items_per_warp = LogicalWarpThreads * ItemsPerThread;
diff --git a/cub/cub/agent/agent_batch_memcpy.cuh b/cub/cub/agent/agent_batch_memcpy.cuh
index 21e4f72c621..f3df08f84d3 100644
--- a/cub/cub/agent/agent_batch_memcpy.cuh
+++ b/cub/cub/agent/agent_batch_memcpy.cuh
@@ -555,18 +555,16 @@ private:
   // TYPE DECLARATIONS
   //---------------------------------------------------------------------
   /// Internal load/store type. For byte-wise memcpy, a single-byte type
-  using AliasT =
-    typename ::cuda::std::conditional<IsMemcpy,
-                                      std::iterator_traits<char*>,
-                                      std::iterator_traits<cub::detail::value_t<InputBufferIt>>>::type::value_type;
+  using AliasT = typename ::cuda::std::
+    conditional_t<IsMemcpy, ::cuda::std::type_identity<char>, lazy_trait<it_value_t, it_value_t<InputBufferIt>>>::type;
 
   /// Types of the input and output buffers
-  using InputBufferT  = cub::detail::value_t<InputBufferIt>;
-  using OutputBufferT = cub::detail::value_t<OutputBufferIt>;
+  using InputBufferT  = it_value_t<InputBufferIt>;
+  using OutputBufferT = it_value_t<OutputBufferIt>;
 
   /// Type that has to be sufficiently large to hold any of the buffers' sizes.
   /// The BufferSizeIteratorT's value type must be convertible to this type.
-  using BufferSizeT = cub::detail::value_t<BufferSizeIteratorT>;
+  using BufferSizeT = it_value_t<BufferSizeIteratorT>;
 
   /// Type used to index into the tile of buffers that this thread block is assigned to.
   using BlockBufferOffsetT = uint16_t;
diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh
index e116b332772..2b5463432da 100644
--- a/cub/cub/agent/agent_histogram.cuh
+++ b/cub/cub/agent/agent_histogram.cuh
@@ -189,7 +189,7 @@ struct AgentHistogram
   //---------------------------------------------------------------------
 
   /// The sample type of the input iterator
-  using SampleT = cub::detail::value_t<SampleIteratorT>;
+  using SampleT = cub::detail::it_value_t<SampleIteratorT>;
 
   /// The pixel type of SampleT
   using PixelT = typename CubVector<SampleT, NUM_CHANNELS>::Type;
diff --git a/cub/cub/agent/agent_merge.cuh b/cub/cub/agent/agent_merge.cuh
index b875fe8cdab..297e800c8b0 100644
--- a/cub/cub/agent/agent_merge.cuh
+++ b/cub/cub/agent/agent_merge.cuh
@@ -61,8 +61,8 @@ struct agent_t
   using policy = Policy;
 
   // key and value type are taken from the first input sequence (consistent with old Thrust behavior)
-  using key_type  = typename ::cuda::std::iterator_traits<KeysIt1>::value_type;
-  using item_type = typename ::cuda::std::iterator_traits<ItemsIt1>::value_type;
+  using key_type  = it_value_t<KeysIt1>;
+  using item_type = it_value_t<ItemsIt1>;
 
   using keys_load_it1  = typename THRUST_NS_QUALIFIER::cuda_cub::core::detail::LoadIterator<Policy, KeysIt1>::type;
   using keys_load_it2  = typename THRUST_NS_QUALIFIER::cuda_cub::core::detail::LoadIterator<Policy, KeysIt2>::type;
diff --git a/cub/cub/agent/agent_reduce.cuh b/cub/cub/agent/agent_reduce.cuh
index c6f7c0a95a4..bd8c8bbf483 100644
--- a/cub/cub/agent/agent_reduce.cuh
+++ b/cub/cub/agent/agent_reduce.cuh
@@ -142,7 +142,7 @@ struct AgentReduce
   //---------------------------------------------------------------------
 
   /// The input value type
-  using InputT = value_t<InputIteratorT>;
+  using InputT = it_value_t<InputIteratorT>;
 
   /// Vector type of InputT for data movement
   using VectorT = typename CubVector<InputT, AgentReducePolicy::VECTOR_LOAD_LENGTH>::Type;
diff --git a/cub/cub/agent/agent_reduce_by_key.cuh b/cub/cub/agent/agent_reduce_by_key.cuh
index 381fa12577f..6852877ec17 100644
--- a/cub/cub/agent/agent_reduce_by_key.cuh
+++ b/cub/cub/agent/agent_reduce_by_key.cuh
@@ -171,13 +171,13 @@ struct AgentReduceByKey
   //---------------------------------------------------------------------
 
   // The input keys type
-  using KeyInputT = value_t<KeysInputIteratorT>;
+  using KeyInputT = it_value_t<KeysInputIteratorT>;
 
   // The output keys type
   using KeyOutputT = non_void_value_t<UniqueOutputIteratorT, KeyInputT>;
 
   // The input values type
-  using ValueInputT = value_t<ValuesInputIteratorT>;
+  using ValueInputT = it_value_t<ValuesInputIteratorT>;
 
   // Tuple type for scanning (pairs accumulated segment-value with
   // segment-index)
diff --git a/cub/cub/agent/agent_rle.cuh b/cub/cub/agent/agent_rle.cuh
index 76fd9f24ed3..fefda393aff 100644
--- a/cub/cub/agent/agent_rle.cuh
+++ b/cub/cub/agent/agent_rle.cuh
@@ -173,7 +173,7 @@ struct AgentRle
   //---------------------------------------------------------------------
 
   /// The input value type
-  using T = cub::detail::value_t<InputIteratorT>;
+  using T = cub::detail::it_value_t<InputIteratorT>;
 
   /// The lengths output value type
   using LengthT = cub::detail::non_void_value_t<LengthsOutputIteratorT, OffsetT>;
diff --git a/cub/cub/agent/agent_scan.cuh b/cub/cub/agent/agent_scan.cuh
index e59fea2ec2e..995ef66811e 100644
--- a/cub/cub/agent/agent_scan.cuh
+++ b/cub/cub/agent/agent_scan.cuh
@@ -157,7 +157,7 @@ struct AgentScan
   //---------------------------------------------------------------------
 
   // The input value type
-  using InputT = cub::detail::value_t<InputIteratorT>;
+  using InputT = cub::detail::it_value_t<InputIteratorT>;
 
   // Tile status descriptor interface type
   using ScanTileStateT = ScanTileState<AccumT>;
diff --git a/cub/cub/agent/agent_scan_by_key.cuh b/cub/cub/agent/agent_scan_by_key.cuh
index a3fcff75464..1ac367742b2 100644
--- a/cub/cub/agent/agent_scan_by_key.cuh
+++ b/cub/cub/agent/agent_scan_by_key.cuh
@@ -145,8 +145,8 @@ struct AgentScanByKey
   // Types and constants
   //---------------------------------------------------------------------
 
-  using KeyT               = value_t<KeysInputIteratorT>;
-  using InputT             = value_t<ValuesInputIteratorT>;
+  using KeyT               = it_value_t<KeysInputIteratorT>;
+  using InputT             = it_value_t<ValuesInputIteratorT>;
   using FlagValuePairT     = KeyValuePair<int, AccumT>;
   using ReduceBySegmentOpT = ScanBySegmentOp<ScanOpT>;
 
diff --git a/cub/cub/agent/agent_select_if.cuh b/cub/cub/agent/agent_select_if.cuh
index 048737382e8..a142e274d2c 100644
--- a/cub/cub/agent/agent_select_if.cuh
+++ b/cub/cub/agent/agent_select_if.cuh
@@ -227,10 +227,10 @@ struct AgentSelectIf
   using MemoryOrderedTileStateT = tile_state_with_memory_order<ScanTileStateT, memory_order>;
 
   // The input value type
-  using InputT = value_t<InputIteratorT>;
+  using InputT = it_value_t<InputIteratorT>;
 
   // The flag value type
-  using FlagT = value_t<FlagsInputIteratorT>;
+  using FlagT = it_value_t<FlagsInputIteratorT>;
 
   // Constants
   enum
diff --git a/cub/cub/agent/agent_three_way_partition.cuh b/cub/cub/agent/agent_three_way_partition.cuh
index 9f2a3afeb62..843167526d4 100644
--- a/cub/cub/agent/agent_three_way_partition.cuh
+++ b/cub/cub/agent/agent_three_way_partition.cuh
@@ -180,7 +180,7 @@ struct AgentThreeWayPartition
   //---------------------------------------------------------------------
 
   // The input value type
-  using InputT = value_t<InputIteratorT>;
+  using InputT = it_value_t<InputIteratorT>;
 
   using AccumPackHelperT = accumulator_pack_t<OffsetT>;
   using AccumPackT       = typename AccumPackHelperT::pack_t;
diff --git a/cub/cub/agent/agent_unique_by_key.cuh b/cub/cub/agent/agent_unique_by_key.cuh
index dfeaaec50a7..1c8a75face9 100644
--- a/cub/cub/agent/agent_unique_by_key.cuh
+++ b/cub/cub/agent/agent_unique_by_key.cuh
@@ -133,8 +133,8 @@ struct AgentUniqueByKey
   //---------------------------------------------------------------------
 
   // The input key and value type
-  using KeyT   = cub::detail::value_t<KeyInputIteratorT>;
-  using ValueT = cub::detail::value_t<ValueInputIteratorT>;
+  using KeyT   = cub::detail::it_value_t<KeyInputIteratorT>;
+  using ValueT = cub::detail::it_value_t<ValueInputIteratorT>;
 
   // Tile status descriptor interface type
   using ScanTileStateT = ScanTileState<OffsetT>;
diff --git a/cub/cub/block/block_load.cuh b/cub/cub/block/block_load.cuh
index 017f0b5ef3f..3bea48e02e0 100644
--- a/cub/cub/block/block_load.cuh
+++ b/cub/cub/block/block_load.cuh
@@ -1250,7 +1250,7 @@ public:
   //! @}  end member group
 };
 
-template <class Policy, class It, class T = cub::detail::value_t<It>>
+template <class Policy, class It, class T = cub::detail::it_value_t<It>>
 struct BlockLoadType
 {
   using type = cub::BlockLoad<T, Policy::BLOCK_THREADS, Policy::ITEMS_PER_THREAD, Policy::LOAD_ALGORITHM>;
diff --git a/cub/cub/block/block_merge_sort.cuh b/cub/cub/block/block_merge_sort.cuh
index 33eb7de6e16..d93b6033084 100644
--- a/cub/cub/block/block_merge_sort.cuh
+++ b/cub/cub/block/block_merge_sort.cuh
@@ -66,8 +66,8 @@ MergePath(KeyIt1 keys1, KeyIt2 keys2, OffsetT keys1_count, OffsetT keys2_count,
   {
     const OffsetT mid = cub::MidPoint<OffsetT>(keys1_begin, keys1_end);
     // pull copies of the keys before calling binary_pred so proxy references are unwrapped
-    const detail::value_t<KeyIt1> key1 = keys1[mid];
-    const detail::value_t<KeyIt2> key2 = keys2[diag - 1 - mid];
+    const detail::it_value_t<KeyIt1> key1 = keys1[mid];
+    const detail::it_value_t<KeyIt2> key2 = keys2[diag - 1 - mid];
     if (binary_pred(key2, key1))
     {
       keys1_end = mid;
diff --git a/cub/cub/block/block_store.cuh b/cub/cub/block/block_store.cuh
index a2cd74fcd90..ea179399b4e 100644
--- a/cub/cub/block/block_store.cuh
+++ b/cub/cub/block/block_store.cuh
@@ -1227,7 +1227,7 @@ public:
 };
 
 #ifndef _CCCL_DOXYGEN_INVOKED // Do not document
-template <class Policy, class It, class T = cub::detail::value_t<It>>
+template <class Policy, class It, class T = cub::detail::it_value_t<It>>
 struct BlockStoreType
 {
   using type = cub::BlockStore<T, Policy::BLOCK_THREADS, Policy::ITEMS_PER_THREAD, Policy::STORE_ALGORITHM>;
diff --git a/cub/cub/device/device_for.cuh b/cub/cub/device/device_for.cuh
index 9359a176905..a0e15dbdd16 100644
--- a/cub/cub/device/device_for.cuh
+++ b/cub/cub/device/device_for.cuh
@@ -44,15 +44,16 @@
 
 #include <thrust/detail/raw_reference_cast.h>
 #include <thrust/distance.h>
-#include <thrust/iterator/iterator_traits.h>
 #include <thrust/system/cuda/detail/core/util.h>
 #include <thrust/type_traits/is_contiguous_iterator.h>
 #include <thrust/type_traits/unwrap_contiguous_iterator.h>
 
+#include <cuda/std/iterator>
+#include <cuda/std/type_traits>
+
 #if __cccl_lib_mdspan
 #  include <cuda/std/__mdspan/extents.h>
 #endif // __cccl_lib_mdspan
-#include <cuda/std/type_traits>
 
 CUB_NAMESPACE_BEGIN
 
@@ -155,7 +156,8 @@ private:
     ContiguousIteratorT first, OffsetT num_items, OpT op, cudaStream_t stream, ::cuda::std::true_type /* vectorize */)
   {
     auto* unwrapped_first = THRUST_NS_QUALIFIER::unwrap_contiguous_iterator(first);
-    using wrapped_op_t = detail::for_each::op_wrapper_vectorized_t<OffsetT, OpT, detail::value_t<ContiguousIteratorT>>;
+    using wrapped_op_t =
+      detail::for_each::op_wrapper_vectorized_t<OffsetT, OpT, detail::it_value_t<ContiguousIteratorT>>;
 
     if (is_aligned<typename wrapped_op_t::vector_t>(unwrapped_first))
     { // Vectorize loads
@@ -594,7 +596,7 @@ private:
     using offset_t = NumItemsT;
     // Disable auto-vectorization for now:
     // constexpr bool use_vectorization =
-    //   detail::for_each::can_regain_copy_freedom<detail::value_t<RandomAccessIteratorT>, OpT>::value
+    //   detail::for_each::can_regain_copy_freedom<detail::it_value_t<RandomAccessIteratorT>, OpT>::value
     //   && THRUST_NS_QUALIFIER::is_contiguous_iterator<RandomAccessIteratorT>::value;
     using use_vectorization_t = ::cuda::std::bool_constant<false>;
     return for_each_n<RandomAccessIteratorT, offset_t, OpT>(first, num_items, op, stream, use_vectorization_t{});
@@ -706,10 +708,8 @@ public:
   {
     CUB_DETAIL_NVTX_RANGE_SCOPE("cub::DeviceFor::ForEach");
 
-    using offset_t = typename THRUST_NS_QUALIFIER::iterator_traits<RandomAccessIteratorT>::difference_type;
-
+    using offset_t       = detail::it_difference_t<RandomAccessIteratorT>;
     const auto num_items = static_cast<offset_t>(THRUST_NS_QUALIFIER::distance(first, last));
-
     return ForEachNNoNVTX(first, num_items, op, stream);
   }
 
@@ -835,7 +835,7 @@ public:
   ForEachCopy(RandomAccessIteratorT first, RandomAccessIteratorT last, OpT op, cudaStream_t stream = {})
   {
     CUB_DETAIL_NVTX_RANGE_SCOPE("cub::DeviceFor::ForEachCopy");
-    using offset_t       = typename THRUST_NS_QUALIFIER::iterator_traits<RandomAccessIteratorT>::difference_type;
+    using offset_t       = detail::it_difference_t<RandomAccessIteratorT>;
     const auto num_items = static_cast<offset_t>(THRUST_NS_QUALIFIER::distance(first, last));
     return ForEachCopyNNoNVTX(first, num_items, op, stream);
   }
diff --git a/cub/cub/device/device_histogram.cuh b/cub/cub/device/device_histogram.cuh
index 29a8721a69a..eacda52673d 100644
--- a/cub/cub/device/device_histogram.cuh
+++ b/cub/cub/device/device_histogram.cuh
@@ -190,7 +190,7 @@ struct DeviceHistogram
     cudaStream_t stream = 0)
   {
     /// The sample value type of the input iterator
-    using SampleT = cub::detail::value_t<SampleIteratorT>;
+    using SampleT = cub::detail::it_value_t<SampleIteratorT>;
     return MultiHistogramEven<1, 1>(
       d_temp_storage,
       temp_storage_bytes,
@@ -509,7 +509,7 @@ struct DeviceHistogram
     cudaStream_t stream = 0)
   {
     /// The sample value type of the input iterator
-    using SampleT = cub::detail::value_t<SampleIteratorT>;
+    using SampleT = cub::detail::it_value_t<SampleIteratorT>;
 
     return MultiHistogramEven<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(
       d_temp_storage,
@@ -700,7 +700,7 @@ struct DeviceHistogram
     CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceHistogram::MultiHistogramEven");
 
     /// The sample value type of the input iterator
-    using SampleT = cub::detail::value_t<SampleIteratorT>;
+    using SampleT = cub::detail::it_value_t<SampleIteratorT>;
     ::cuda::std::bool_constant<sizeof(SampleT) == 1> is_byte_sample;
 
     if constexpr (sizeof(OffsetT) > sizeof(int))
@@ -850,7 +850,7 @@ struct DeviceHistogram
     cudaStream_t stream = 0)
   {
     /// The sample value type of the input iterator
-    using SampleT = cub::detail::value_t<SampleIteratorT>;
+    using SampleT = cub::detail::it_value_t<SampleIteratorT>;
     return MultiHistogramRange<1, 1>(
       d_temp_storage,
       temp_storage_bytes,
@@ -1145,7 +1145,7 @@ struct DeviceHistogram
     cudaStream_t stream = 0)
   {
     /// The sample value type of the input iterator
-    using SampleT = cub::detail::value_t<SampleIteratorT>;
+    using SampleT = cub::detail::it_value_t<SampleIteratorT>;
 
     return MultiHistogramRange<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(
       d_temp_storage,
@@ -1326,7 +1326,7 @@ struct DeviceHistogram
     CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceHistogram::MultiHistogramRange");
 
     /// The sample value type of the input iterator
-    using SampleT = cub::detail::value_t<SampleIteratorT>;
+    using SampleT = cub::detail::it_value_t<SampleIteratorT>;
     ::cuda::std::bool_constant<sizeof(SampleT) == 1> is_byte_sample;
 
     if constexpr (sizeof(OffsetT) > sizeof(int))
diff --git a/cub/cub/device/device_memcpy.cuh b/cub/cub/device/device_memcpy.cuh
index 997dfecd4c3..68dcc806bc5 100644
--- a/cub/cub/device/device_memcpy.cuh
+++ b/cub/cub/device/device_memcpy.cuh
@@ -175,10 +175,10 @@ struct DeviceMemcpy
     cudaStream_t stream = 0)
   {
     CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceMemcpy::Batched");
-    static_assert(::cuda::std::is_pointer_v<cub::detail::value_t<InputBufferIt>>,
+    static_assert(::cuda::std::is_pointer_v<cub::detail::it_value_t<InputBufferIt>>,
                   "DeviceMemcpy::Batched only supports copying of memory buffers."
                   "Please consider using DeviceCopy::Batched instead.");
-    static_assert(::cuda::std::is_pointer_v<cub::detail::value_t<OutputBufferIt>>,
+    static_assert(::cuda::std::is_pointer_v<cub::detail::it_value_t<OutputBufferIt>>,
                   "DeviceMemcpy::Batched only supports copying of memory buffers."
                   "Please consider using DeviceCopy::Batched instead.");
 
diff --git a/cub/cub/device/device_reduce.cuh b/cub/cub/device/device_reduce.cuh
index 93eb2024f67..b286afa31df 100644
--- a/cub/cub/device/device_reduce.cuh
+++ b/cub/cub/device/device_reduce.cuh
@@ -318,7 +318,7 @@ struct DeviceReduce
     using OffsetT = detail::choose_offset_t<NumItemsT>;
 
     // The output value type
-    using OutputT = cub::detail::non_void_value_t<OutputIteratorT, cub::detail::value_t<InputIteratorT>>;
+    using OutputT = cub::detail::non_void_value_t<OutputIteratorT, cub::detail::it_value_t<InputIteratorT>>;
 
     using InitT = OutputT;
 
@@ -424,7 +424,7 @@ struct DeviceReduce
     using OffsetT = detail::choose_offset_t<NumItemsT>;
 
     // The input value type
-    using InputT = cub::detail::value_t<InputIteratorT>;
+    using InputT = cub::detail::it_value_t<InputIteratorT>;
 
     using InitT = InputT;
 
@@ -539,7 +539,7 @@ struct DeviceReduce
     CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceReduce::ArgMin");
 
     // The input type
-    using InputValueT = cub::detail::value_t<InputIteratorT>;
+    using InputValueT = cub::detail::it_value_t<InputIteratorT>;
 
     // Offset type used within the kernel and to index within one partition
     using PerPartitionOffsetT = int;
@@ -673,7 +673,7 @@ struct DeviceReduce
     using OffsetT = int;
 
     // The input type
-    using InputValueT = cub::detail::value_t<InputIteratorT>;
+    using InputValueT = cub::detail::it_value_t<InputIteratorT>;
 
     // The output tuple type
     using OutputTupleT = cub::detail::non_void_value_t<OutputIteratorT, KeyValuePair<OffsetT, InputValueT>>;
@@ -785,7 +785,7 @@ struct DeviceReduce
     using OffsetT = detail::choose_offset_t<NumItemsT>;
 
     // The input value type
-    using InputT = cub::detail::value_t<InputIteratorT>;
+    using InputT = cub::detail::it_value_t<InputIteratorT>;
 
     using InitT = InputT;
 
@@ -900,7 +900,7 @@ struct DeviceReduce
     CUB_DETAIL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceReduce::ArgMax");
 
     // The input type
-    using InputValueT = cub::detail::value_t<InputIteratorT>;
+    using InputValueT = cub::detail::it_value_t<InputIteratorT>;
 
     // Offset type used within the kernel and to index within one partition
     using PerPartitionOffsetT = int;
@@ -1038,7 +1038,7 @@ struct DeviceReduce
     using OffsetT = int;
 
     // The input type
-    using InputValueT = cub::detail::value_t<InputIteratorT>;
+    using InputValueT = cub::detail::it_value_t<InputIteratorT>;
 
     // The output tuple type
     using OutputTupleT = cub::detail::non_void_value_t<OutputIteratorT, KeyValuePair<OffsetT, InputValueT>>;
diff --git a/cub/cub/device/device_run_length_encode.cuh b/cub/cub/device/device_run_length_encode.cuh
index 21ed7bd77ec..b308f4c2b65 100644
--- a/cub/cub/device/device_run_length_encode.cuh
+++ b/cub/cub/device/device_run_length_encode.cuh
@@ -205,7 +205,7 @@ struct DeviceRunLengthEncode
 
     using accum_t = ::cuda::std::__accumulator_t<reduction_op, length_t, length_t>;
 
-    using key_t = cub::detail::non_void_value_t<UniqueOutputIteratorT, cub::detail::value_t<InputIteratorT>>;
+    using key_t = cub::detail::non_void_value_t<UniqueOutputIteratorT, cub::detail::it_value_t<InputIteratorT>>;
 
     using policy_t = detail::rle::encode::policy_hub<accum_t, key_t>;
 
diff --git a/cub/cub/device/device_scan.cuh b/cub/cub/device/device_scan.cuh
index 74f351f1060..84514d3ff4d 100644
--- a/cub/cub/device/device_scan.cuh
+++ b/cub/cub/device/device_scan.cuh
@@ -191,7 +191,7 @@ struct DeviceScan
 
     // Unsigned integer type for global offsets
     using OffsetT = detail::choose_offset_t<NumItemsT>;
-    using InitT   = cub::detail::value_t<InputIteratorT>;
+    using InitT   = cub::detail::it_value_t<InputIteratorT>;
 
     // Initial value
     InitT init_value{};
@@ -1156,7 +1156,7 @@ struct DeviceScan
 
     // Unsigned integer type for global offsets
     using OffsetT = detail::choose_offset_t<NumItemsT>;
-    using AccumT  = ::cuda::std::__accumulator_t<ScanOpT, cub::detail::value_t<InputIteratorT>, InitValueT>;
+    using AccumT  = ::cuda::std::__accumulator_t<ScanOpT, cub::detail::it_value_t<InputIteratorT>, InitValueT>;
 
     return DispatchScan<
       InputIteratorT,
@@ -1390,7 +1390,7 @@ struct DeviceScan
 
     // Unsigned integer type for global offsets
     using OffsetT = detail::choose_offset_t<NumItemsT>;
-    using InitT   = cub::detail::value_t<ValuesInputIteratorT>;
+    using InitT   = cub::detail::it_value_t<ValuesInputIteratorT>;
 
     // Initial value
     InitT init_value{};
diff --git a/cub/cub/device/device_segmented_reduce.cuh b/cub/cub/device/device_segmented_reduce.cuh
index 551c487a847..35cb0e81132 100644
--- a/cub/cub/device/device_segmented_reduce.cuh
+++ b/cub/cub/device/device_segmented_reduce.cuh
@@ -366,7 +366,7 @@ public:
     using OffsetT = detail::common_iterator_value_t<BeginOffsetIteratorT, EndOffsetIteratorT>;
 
     // The output value type
-    using OutputT               = cub::detail::non_void_value_t<OutputIteratorT, cub::detail::value_t<InputIteratorT>>;
+    using OutputT = cub::detail::non_void_value_t<OutputIteratorT, cub::detail::it_value_t<InputIteratorT>>;
     using integral_offset_check = ::cuda::std::is_integral<OffsetT>;
 
     static_assert(integral_offset_check::value, "Offset iterator value type should be integral.");
@@ -489,7 +489,7 @@ public:
     using OffsetT = detail::common_iterator_value_t<BeginOffsetIteratorT, EndOffsetIteratorT>;
 
     // The input value type
-    using InputT                = cub::detail::value_t<InputIteratorT>;
+    using InputT                = cub::detail::it_value_t<InputIteratorT>;
     using integral_offset_check = ::cuda::std::is_integral<OffsetT>;
 
     static_assert(integral_offset_check::value, "Offset iterator value type should be integral.");
@@ -618,7 +618,7 @@ public:
     using OffsetT = int; // detail::common_iterator_value_t<BeginOffsetIteratorT, EndOffsetIteratorT>;
 
     // The input type
-    using InputValueT = cub::detail::value_t<InputIteratorT>;
+    using InputValueT = cub::detail::it_value_t<InputIteratorT>;
 
     // The output tuple type
     using OutputTupleT = cub::detail::non_void_value_t<OutputIteratorT, KeyValuePair<OffsetT, InputValueT>>;
@@ -755,7 +755,7 @@ public:
     using OffsetT = detail::common_iterator_value_t<BeginOffsetIteratorT, EndOffsetIteratorT>;
 
     // The input value type
-    using InputT = cub::detail::value_t<InputIteratorT>;
+    using InputT = cub::detail::it_value_t<InputIteratorT>;
 
     using integral_offset_check = ::cuda::std::is_integral<OffsetT>;
     static_assert(integral_offset_check::value, "Offset iterator value type should be integral.");
@@ -882,7 +882,7 @@ public:
     using OffsetT = int; // detail::common_iterator_value_t<BeginOffsetIteratorT, EndOffsetIteratorT>;
 
     // The input type
-    using InputValueT = cub::detail::value_t<InputIteratorT>;
+    using InputValueT = cub::detail::it_value_t<InputIteratorT>;
 
     // The output tuple type
     using OutputTupleT = cub::detail::non_void_value_t<OutputIteratorT, KeyValuePair<OffsetT, InputValueT>>;
diff --git a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh
index 60f6585cf52..f74fc17ab94 100644
--- a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh
+++ b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh
@@ -122,7 +122,7 @@ template <typename InputIteratorT,
           typename PolicyHub = detail::adjacent_difference::policy_hub<InputIteratorT, AliasOpt == MayAlias::Yes>>
 struct DispatchAdjacentDifference
 {
-  using InputT = typename std::iterator_traits<InputIteratorT>::value_type;
+  using InputT = detail::it_value_t<InputIteratorT>;
 
   void* d_temp_storage;
   size_t& temp_storage_bytes;
diff --git a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh
index f24c14df86b..c47fc8ae016 100644
--- a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh
+++ b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh
@@ -56,6 +56,7 @@
 
 #include <cuda/std/__algorithm/max.h>
 #include <cuda/std/__algorithm/min.h>
+#include <cuda/std/iterator>
 #include <cuda/std/type_traits>
 
 #include <cstdint>
@@ -110,14 +111,14 @@ __launch_bounds__(int(ChainedPolicyT::ActivePolicy::AgentLargeBufferPolicyT::BLO
 {
   using StatusWord    = typename TileT::StatusWord;
   using ActivePolicyT = typename ChainedPolicyT::ActivePolicy::AgentLargeBufferPolicyT;
-  using BufferSizeT   = value_t<BufferSizeIteratorT>;
+  using BufferSizeT   = it_value_t<BufferSizeIteratorT>;
   /// Internal load/store type. For byte-wise memcpy, a single-byte type
   using AliasT = typename ::cuda::std::conditional_t<MemcpyOpt == CopyAlg::Memcpy,
-                                                     std::iterator_traits<char*>,
-                                                     std::iterator_traits<value_t<InputBufferIt>>>::value_type;
+                                                     ::cuda::std::type_identity<char>,
+                                                     lazy_trait<it_value_t, it_value_t<InputBufferIt>>>::type;
   /// Types of the input and output buffers
-  using InputBufferT  = value_t<InputBufferIt>;
-  using OutputBufferT = value_t<OutputBufferIt>;
+  using InputBufferT  = it_value_t<InputBufferIt>;
+  using OutputBufferT = it_value_t<OutputBufferIt>;
 
   constexpr uint32_t BLOCK_THREADS    = ActivePolicyT::BLOCK_THREADS;
   constexpr uint32_t ITEMS_PER_THREAD = ActivePolicyT::BYTES_PER_THREAD;
@@ -240,7 +241,7 @@ __launch_bounds__(int(ChainedPolicyT::ActivePolicy::AgentSmallBufferPolicyT::BLO
     BLevBlockOffsetTileState blev_block_scan_state)
 {
   // Internal type used for storing a buffer's size
-  using BufferSizeT = value_t<BufferSizeIteratorT>;
+  using BufferSizeT = it_value_t<BufferSizeIteratorT>;
 
   // Alias the correct tuning policy for the current compilation pass' architecture
   using AgentBatchMemcpyPolicyT = typename ChainedPolicyT::ActivePolicy::AgentSmallBufferPolicyT;
@@ -313,7 +314,7 @@ struct DispatchBatchMemcpy
   using BufferTileOffsetScanStateT = typename cub::ScanTileState<BlockOffsetT>;
 
   // Internal type used to keep track of a buffer's size
-  using BufferSizeT = cub::detail::value_t<BufferSizeIteratorT>;
+  using BufferSizeT = cub::detail::it_value_t<BufferSizeIteratorT>;
 
   //------------------------------------------------------------------------------
   // Member Variables
@@ -397,9 +398,9 @@ struct DispatchBatchMemcpy
     BlockOffsetT num_tiles = ::cuda::ceil_div(num_buffers, TILE_SIZE);
 
     using BlevBufferSrcsOutT =
-      ::cuda::std::_If<MemcpyOpt == CopyAlg::Memcpy, const void*, cub::detail::value_t<InputBufferIt>>;
+      ::cuda::std::_If<MemcpyOpt == CopyAlg::Memcpy, const void*, cub::detail::it_value_t<InputBufferIt>>;
     using BlevBufferDstOutT =
-      ::cuda::std::_If<MemcpyOpt == CopyAlg::Memcpy, void*, cub::detail::value_t<OutputBufferIt>>;
+      ::cuda::std::_If<MemcpyOpt == CopyAlg::Memcpy, void*, cub::detail::it_value_t<OutputBufferIt>>;
     using BlevBufferSrcsOutItT        = BlevBufferSrcsOutT*;
     using BlevBufferDstsOutItT        = BlevBufferDstOutT*;
     using BlevBufferSizesOutItT       = BufferSizeT*;
diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh
index f63af0fa2f4..d5a894c3310 100644
--- a/cub/cub/device/dispatch/dispatch_histogram.cuh
+++ b/cub/cub/device/dispatch/dispatch_histogram.cuh
@@ -569,7 +569,7 @@ public:
   //---------------------------------------------------------------------
 
   /// The sample value type of the input iterator
-  using SampleT = cub::detail::value_t<SampleIteratorT>;
+  using SampleT = cub::detail::it_value_t<SampleIteratorT>;
 
   enum
   {
@@ -924,7 +924,7 @@ public:
     // Should we call DispatchHistogram<....., PolicyHub=void> in DeviceHistogram?
     static constexpr bool isEven = 0;
     using fallback_policy_hub    = detail::histogram::
-      policy_hub<detail::value_t<SampleIteratorT>, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, isEven>;
+      policy_hub<detail::it_value_t<SampleIteratorT>, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, isEven>;
 
     using MaxPolicyT =
       typename ::cuda::std::_If<::cuda::std::is_void_v<PolicyHub>, fallback_policy_hub, PolicyHub>::MaxPolicy;
@@ -1100,7 +1100,7 @@ public:
   {
     static constexpr bool isEven = 0;
     using fallback_policy_hub    = detail::histogram::
-      policy_hub<detail::value_t<SampleIteratorT>, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, isEven>;
+      policy_hub<detail::it_value_t<SampleIteratorT>, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, isEven>;
 
     using MaxPolicyT =
       typename ::cuda::std::_If<::cuda::std::is_void_v<PolicyHub>, fallback_policy_hub, PolicyHub>::MaxPolicy;
@@ -1240,7 +1240,7 @@ public:
   {
     static constexpr bool isEven = 1;
     using fallback_policy_hub    = detail::histogram::
-      policy_hub<detail::value_t<SampleIteratorT>, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, isEven>;
+      policy_hub<detail::it_value_t<SampleIteratorT>, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, isEven>;
 
     using MaxPolicyT =
       typename ::cuda::std::_If<::cuda::std::is_void_v<PolicyHub>, fallback_policy_hub, PolicyHub>::MaxPolicy;
@@ -1431,7 +1431,7 @@ public:
   {
     static constexpr bool isEven = 1;
     using fallback_policy_hub    = detail::histogram::
-      policy_hub<detail::value_t<SampleIteratorT>, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, isEven>;
+      policy_hub<detail::it_value_t<SampleIteratorT>, CounterT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, isEven>;
 
     using MaxPolicyT =
       typename ::cuda::std::_If<::cuda::std::is_void_v<PolicyHub>, fallback_policy_hub, PolicyHub>::MaxPolicy;
diff --git a/cub/cub/device/dispatch/dispatch_merge.cuh b/cub/cub/device/dispatch/dispatch_merge.cuh
index f9b40d053a5..af65a150d99 100644
--- a/cub/cub/device/dispatch/dispatch_merge.cuh
+++ b/cub/cub/device/dispatch/dispatch_merge.cuh
@@ -119,7 +119,7 @@ __launch_bounds__(
     vsmem_t global_temp_storage)
 {
   // the merge agent loads keys into a local array of KeyIt1::value_type, on which the comparisons are performed
-  using key_t = value_t<KeyIt1>;
+  using key_t = it_value_t<KeyIt1>;
   static_assert(::cuda::std::__invokable<CompareOp, key_t, key_t>::value,
                 "Comparison operator cannot compare two keys");
   static_assert(::cuda::std::is_convertible_v<typename ::cuda::std::__invoke_of<CompareOp, key_t, key_t>::type, bool>,
@@ -164,7 +164,7 @@ template <typename KeyIt1,
           typename ValueIt3,
           typename Offset,
           typename CompareOp,
-          typename PolicyHub = detail::merge::policy_hub<value_t<KeyIt1>, value_t<ValueIt1>>>
+          typename PolicyHub = detail::merge::policy_hub<it_value_t<KeyIt1>, it_value_t<ValueIt1>>>
 struct dispatch_t
 {
   void* d_temp_storage;
diff --git a/cub/cub/device/dispatch/dispatch_merge_sort.cuh b/cub/cub/device/dispatch/dispatch_merge_sort.cuh
index 330a30dd940..6cd0c75a3fa 100644
--- a/cub/cub/device/dispatch/dispatch_merge_sort.cuh
+++ b/cub/cub/device/dispatch/dispatch_merge_sort.cuh
@@ -64,8 +64,8 @@ template <typename MaxPolicyT,
           typename CompareOpT>
 struct DeviceMergeSortKernelSource
 {
-  using KeyT   = cub::detail::value_t<KeyIteratorT>;
-  using ValueT = cub::detail::value_t<ValueIteratorT>;
+  using KeyT   = cub::detail::it_value_t<KeyIteratorT>;
+  using ValueT = cub::detail::it_value_t<ValueIteratorT>;
 
   CUB_DEFINE_KERNEL_GETTER(
     MergeSortBlockSortKernel,
@@ -127,10 +127,10 @@ template <
        ValueIteratorT,
        OffsetT,
        CompareOpT,
-       cub::detail::value_t<KeyIteratorT>,
-       cub::detail::value_t<ValueIteratorT>>,
-  typename KeyT   = cub::detail::value_t<KeyIteratorT>,
-  typename ValueT = cub::detail::value_t<ValueIteratorT>>
+       cub::detail::it_value_t<KeyIteratorT>,
+       cub::detail::it_value_t<ValueIteratorT>>,
+  typename KeyT   = cub::detail::it_value_t<KeyIteratorT>,
+  typename ValueT = cub::detail::it_value_t<ValueIteratorT>>
 struct DispatchMergeSort
 {
   /// Whether or not there are values to be trucked along with keys
diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh
index bc9faead2e3..8ac66edb6b0 100644
--- a/cub/cub/device/dispatch/dispatch_reduce.cuh
+++ b/cub/cub/device/dispatch/dispatch_reduce.cuh
@@ -137,8 +137,8 @@ template <typename InputIteratorT,
           typename OutputIteratorT,
           typename OffsetT,
           typename ReductionOpT,
-          typename InitT  = cub::detail::non_void_value_t<OutputIteratorT, cub::detail::value_t<InputIteratorT>>,
-          typename AccumT = ::cuda::std::__accumulator_t<ReductionOpT, cub::detail::value_t<InputIteratorT>, InitT>,
+          typename InitT  = cub::detail::non_void_value_t<OutputIteratorT, cub::detail::it_value_t<InputIteratorT>>,
+          typename AccumT = ::cuda::std::__accumulator_t<ReductionOpT, cub::detail::it_value_t<InputIteratorT>, InitT>,
           typename TransformOpT = ::cuda::std::__identity,
           typename PolicyHub    = detail::reduce::policy_hub<AccumT, OffsetT, ReductionOpT>,
           typename KernelSource = detail::reduce::DeviceReduceKernelSource<
@@ -573,8 +573,10 @@ template <
   typename ReductionOpT,
   typename TransformOpT,
   typename InitT,
-  typename AccumT = ::cuda::std::
-    __accumulator_t<ReductionOpT, cub::detail::invoke_result_t<TransformOpT, cub::detail::value_t<InputIteratorT>>, InitT>,
+  typename AccumT =
+    ::cuda::std::__accumulator_t<ReductionOpT,
+                                 cub::detail::invoke_result_t<TransformOpT, cub::detail::it_value_t<InputIteratorT>>,
+                                 InitT>,
   typename PolicyHub    = detail::reduce::policy_hub<AccumT, OffsetT, ReductionOpT>,
   typename KernelSource = detail::reduce::DeviceReduceKernelSource<
     typename PolicyHub::MaxPolicy,
@@ -671,9 +673,9 @@ template <typename InputIteratorT,
           typename EndOffsetIteratorT,
           typename OffsetT,
           typename ReductionOpT,
-          typename InitT     = cub::detail::non_void_value_t<OutputIteratorT, cub::detail::value_t<InputIteratorT>>,
-          typename AccumT    = ::cuda::std::__accumulator_t<ReductionOpT, cub::detail::value_t<InputIteratorT>, InitT>,
-          typename PolicyHub = detail::reduce::policy_hub<AccumT, OffsetT, ReductionOpT>,
+          typename InitT  = cub::detail::non_void_value_t<OutputIteratorT, cub::detail::it_value_t<InputIteratorT>>,
+          typename AccumT = ::cuda::std::__accumulator_t<ReductionOpT, cub::detail::it_value_t<InputIteratorT>, InitT>,
+          typename PolicyHub    = detail::reduce::policy_hub<AccumT, OffsetT, ReductionOpT>,
           typename KernelSource = detail::reduce::DeviceSegmentedReduceKernelSource<
             typename PolicyHub::MaxPolicy,
             InputIteratorT,
diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
index 439dbd48634..aa95fdd1342 100644
--- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
+++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
@@ -226,12 +226,12 @@ template <typename KeysInputIteratorT,
           typename ReductionOpT,
           typename OffsetT,
           typename AccumT    = ::cuda::std::__accumulator_t<ReductionOpT,
-                                                            cub::detail::value_t<ValuesInputIteratorT>,
-                                                            cub::detail::value_t<ValuesInputIteratorT>>,
+                                                            cub::detail::it_value_t<ValuesInputIteratorT>,
+                                                            cub::detail::it_value_t<ValuesInputIteratorT>>,
           typename PolicyHub = detail::reduce_by_key::policy_hub<
             ReductionOpT,
             AccumT,
-            cub::detail::non_void_value_t<UniqueOutputIteratorT, cub::detail::value_t<KeysInputIteratorT>>>>
+            cub::detail::non_void_value_t<UniqueOutputIteratorT, cub::detail::it_value_t<KeysInputIteratorT>>>>
 struct DispatchReduceByKey
 {
   //-------------------------------------------------------------------------
@@ -239,7 +239,7 @@ struct DispatchReduceByKey
   //-------------------------------------------------------------------------
 
   // The input values type
-  using ValueInputT = cub::detail::value_t<ValuesInputIteratorT>;
+  using ValueInputT = cub::detail::it_value_t<ValuesInputIteratorT>;
 
   static constexpr int INIT_KERNEL_THREADS = 128;
 
diff --git a/cub/cub/device/dispatch/dispatch_rle.cuh b/cub/cub/device/dispatch/dispatch_rle.cuh
index 75e15980427..3ba14bd124f 100644
--- a/cub/cub/device/dispatch/dispatch_rle.cuh
+++ b/cub/cub/device/dispatch/dispatch_rle.cuh
@@ -194,7 +194,7 @@ template <typename InputIteratorT,
           typename OffsetT,
           typename PolicyHub =
             detail::rle::non_trivial_runs::policy_hub<cub::detail::non_void_value_t<LengthsOutputIteratorT, OffsetT>,
-                                                      cub::detail::value_t<InputIteratorT>>>
+                                                      cub::detail::it_value_t<InputIteratorT>>>
 struct DeviceRleDispatch
 {
   /******************************************************************************
diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh
index 85834b10a35..fe0bc307d4f 100644
--- a/cub/cub/device/dispatch/dispatch_scan.cuh
+++ b/cub/cub/device/dispatch/dispatch_scan.cuh
@@ -145,13 +145,13 @@ template <
   typename InitValueT,
   typename OffsetT,
   typename AccumT                 = ::cuda::std::__accumulator_t<ScanOpT,
-                                                                 cub::detail::value_t<InputIteratorT>,
+                                                                 cub::detail::it_value_t<InputIteratorT>,
                                                                  ::cuda::std::_If<::cuda::std::is_same_v<InitValueT, NullType>,
-                                                                                  cub::detail::value_t<InputIteratorT>,
+                                                                                  cub::detail::it_value_t<InputIteratorT>,
                                                                                   typename InitValueT::value_type>>,
   ForceInclusive EnforceInclusive = ForceInclusive::No,
-  typename PolicyHub =
-    detail::scan::policy_hub<detail::value_t<InputIteratorT>, detail::value_t<OutputIteratorT>, AccumT, OffsetT, ScanOpT>,
+  typename PolicyHub              = detail::scan::
+    policy_hub<detail::it_value_t<InputIteratorT>, detail::it_value_t<OutputIteratorT>, AccumT, OffsetT, ScanOpT>,
   typename KernelSource = detail::scan::DeviceScanKernelSource<
     typename PolicyHub::MaxPolicy,
     InputIteratorT,
diff --git a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh
index 3290744fab5..c8db2e84995 100644
--- a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh
+++ b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh
@@ -136,7 +136,7 @@ template <typename ChainedPolicyT,
           typename InitValueT,
           typename OffsetT,
           typename AccumT,
-          typename KeyT = cub::detail::value_t<KeysInputIteratorT>>
+          typename KeyT = cub::detail::it_value_t<KeysInputIteratorT>>
 __launch_bounds__(int(ChainedPolicyT::ActivePolicy::ScanByKeyPolicyT::BLOCK_THREADS))
   CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceScanByKeyKernel(
     KeysInputIteratorT d_keys_in,
@@ -176,7 +176,7 @@ template <typename ScanTileStateT, typename KeysInputIteratorT, typename OffsetT
 CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceScanByKeyInitKernel(
   ScanTileStateT tile_state,
   KeysInputIteratorT d_keys_in,
-  cub::detail::value_t<KeysInputIteratorT>* d_keys_prev_in,
+  cub::detail::it_value_t<KeysInputIteratorT>* d_keys_prev_in,
   OffsetT items_per_tile,
   int num_tiles)
 {
@@ -233,10 +233,11 @@ template <
   typename OffsetT,
   typename AccumT = ::cuda::std::__accumulator_t<
     ScanOpT,
-    cub::detail::value_t<ValuesInputIteratorT>,
-    ::cuda::std::_If<::cuda::std::is_same_v<InitValueT, NullType>, cub::detail::value_t<ValuesInputIteratorT>, InitValueT>>,
+    cub::detail::it_value_t<ValuesInputIteratorT>,
+    ::cuda::std::
+      _If<::cuda::std::is_same_v<InitValueT, NullType>, cub::detail::it_value_t<ValuesInputIteratorT>, InitValueT>>,
   typename PolicyHub =
-    detail::scan_by_key::policy_hub<KeysInputIteratorT, AccumT, cub::detail::value_t<ValuesInputIteratorT>, ScanOpT>>
+    detail::scan_by_key::policy_hub<KeysInputIteratorT, AccumT, cub::detail::it_value_t<ValuesInputIteratorT>, ScanOpT>>
 struct DispatchScanByKey
 {
   //---------------------------------------------------------------------
@@ -246,10 +247,10 @@ struct DispatchScanByKey
   static constexpr int INIT_KERNEL_THREADS = 128;
 
   // The input key type
-  using KeyT = cub::detail::value_t<KeysInputIteratorT>;
+  using KeyT = cub::detail::it_value_t<KeysInputIteratorT>;
 
   // The input value type
-  using InputT = cub::detail::value_t<ValuesInputIteratorT>;
+  using InputT = cub::detail::it_value_t<ValuesInputIteratorT>;
 
   // Tile state used for the decoupled look-back
   using ScanByKeyTileStateT = ReduceByKeyScanTileState<AccumT, int>;
diff --git a/cub/cub/device/dispatch/dispatch_select_if.cuh b/cub/cub/device/dispatch/dispatch_select_if.cuh
index eb226ae754e..561e34a1fb8 100644
--- a/cub/cub/device/dispatch/dispatch_select_if.cuh
+++ b/cub/cub/device/dispatch/dispatch_select_if.cuh
@@ -421,8 +421,8 @@ template <
   typename OffsetT,
   SelectImpl SelectionOpt,
   typename PolicyHub = detail::select::policy_hub<
-    detail::value_t<InputIteratorT>,
-    detail::value_t<FlagsInputIteratorT>,
+    detail::it_value_t<InputIteratorT>,
+    detail::it_value_t<FlagsInputIteratorT>,
     // if/flagged/unique only have a single code path for different offset types, partition has different code paths
     ::cuda::std::conditional_t<SelectionOpt == SelectImpl::Partition, OffsetT, detail::select::per_partition_offset_t>,
     detail::select::is_partition_distinct_output_t<SelectedOutputIteratorT>::value,
diff --git a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh
index 28b3a06cbd4..b560de56dc7 100644
--- a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh
+++ b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh
@@ -241,7 +241,7 @@ template <typename InputIteratorT,
           typename SelectSecondPartOp,
           typename OffsetT,
           typename PolicyHub = detail::three_way_partition::
-            policy_hub<cub::detail::value_t<InputIteratorT>, detail::three_way_partition::per_partition_offset_t>>
+            policy_hub<cub::detail::it_value_t<InputIteratorT>, detail::three_way_partition::per_partition_offset_t>>
 struct DispatchThreeWayPartitionIf
 {
   /*****************************************************************************
diff --git a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh
index e30c33a546d..9120319c49d 100644
--- a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh
+++ b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh
@@ -79,15 +79,16 @@ CUB_NAMESPACE_BEGIN
  * @tparam OffsetT
  *   Signed integer type for global offsets
  */
-template <typename KeyInputIteratorT,
-          typename ValueInputIteratorT,
-          typename KeyOutputIteratorT,
-          typename ValueOutputIteratorT,
-          typename NumSelectedIteratorT,
-          typename EqualityOpT,
-          typename OffsetT,
-          typename PolicyHub =
-            detail::unique_by_key::policy_hub<detail::value_t<KeyInputIteratorT>, detail::value_t<ValueInputIteratorT>>>
+template <
+  typename KeyInputIteratorT,
+  typename ValueInputIteratorT,
+  typename KeyOutputIteratorT,
+  typename ValueOutputIteratorT,
+  typename NumSelectedIteratorT,
+  typename EqualityOpT,
+  typename OffsetT,
+  typename PolicyHub =
+    detail::unique_by_key::policy_hub<detail::it_value_t<KeyInputIteratorT>, detail::it_value_t<ValueInputIteratorT>>>
 struct DispatchUniqueByKey
 {
   /******************************************************************************
@@ -100,8 +101,8 @@ struct DispatchUniqueByKey
   };
 
   // The input key and value type
-  using KeyT   = typename std::iterator_traits<KeyInputIteratorT>::value_type;
-  using ValueT = typename std::iterator_traits<ValueInputIteratorT>::value_type;
+  using KeyT   = detail::it_value_t<KeyInputIteratorT>;
+  using ValueT = detail::it_value_t<ValueInputIteratorT>;
 
   // Tile status descriptor interface type
   using ScanTileStateT = ScanTileState<OffsetT>;
diff --git a/cub/cub/device/dispatch/kernels/transform.cuh b/cub/cub/device/dispatch/kernels/transform.cuh
index 51193346036..42ba275c07d 100644
--- a/cub/cub/device/dispatch/kernels/transform.cuh
+++ b/cub/cub/device/dispatch/kernels/transform.cuh
@@ -59,7 +59,7 @@ _CCCL_DEVICE _CCCL_FORCEINLINE void prefetch_tile(It begin, int tile_size)
   {
     constexpr int prefetch_byte_stride = 128; // TODO(bgruber): should correspond to cache line size. Does this need to
                                               // be architecture dependent?
-    const int tile_size_bytes = tile_size * sizeof(value_t<It>);
+    const int tile_size_bytes = tile_size * sizeof(it_value_t<It>);
     // prefetch does not stall and unrolling just generates a lot of unnecessary computations and predicate handling
 #pragma unroll 1
     for (int offset = threadIdx.x * prefetch_byte_stride; offset < tile_size_bytes;
@@ -361,7 +361,7 @@ _CCCL_DEVICE void transform_kernel_impl(
 template <typename It>
 union kernel_arg
 {
-  aligned_base_ptr<value_t<It>> aligned_ptr; // first member is trivial
+  aligned_base_ptr<it_value_t<It>> aligned_ptr; // first member is trivial
   It iterator; // may not be trivially [default|copy]-constructible
 
   static_assert(::cuda::std::is_trivial_v<decltype(aligned_ptr)>, "");
diff --git a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
index b8d0a7557bd..b2fc266bfc9 100644
--- a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh
@@ -50,7 +50,7 @@ namespace adjacent_difference
 template <typename InputIteratorT, bool MayAlias>
 struct policy_hub
 {
-  using ValueT = typename std::iterator_traits<InputIteratorT>::value_type;
+  using ValueT = it_value_t<InputIteratorT>;
 
   struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
   {
diff --git a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
index 2c93b1b1147..8064337ab28 100644
--- a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
@@ -73,7 +73,7 @@ CUB_RUNTIME_FUNCTION MergeSortPolicyWrapper<PolicyT> MakeMergeSortPolicyWrapper(
 template <typename KeyIteratorT>
 struct policy_hub
 {
-  using KeyT = value_t<KeyIteratorT>;
+  using KeyT = it_value_t<KeyIteratorT>;
 
   struct Policy500 : ChainedPolicy<500, Policy500, Policy500>
   {
diff --git a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
index 3723f1d84f6..4695ce58d66 100644
--- a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
@@ -912,7 +912,7 @@ struct sm100_tuning<KeyT, ValueT, primitive_op::yes, key_size::_8, val_size::_8,
 template <typename KeysInputIteratorT, typename AccumT, typename ValueT, typename ScanOpT>
 struct policy_hub
 {
-  using key_t                               = value_t<KeysInputIteratorT>;
+  using key_t                               = it_value_t<KeysInputIteratorT>;
   static constexpr int max_input_bytes      = static_cast<int>((::cuda::std::max)(sizeof(key_t), sizeof(AccumT)));
   static constexpr int combined_input_bytes = static_cast<int>(sizeof(key_t) + sizeof(AccumT));
 
diff --git a/cub/cub/device/dispatch/tuning/tuning_transform.cuh b/cub/cub/device/dispatch/tuning/tuning_transform.cuh
index 29c1a6d1a9b..ddd0482a69c 100644
--- a/cub/cub/device/dispatch/tuning/tuning_transform.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_transform.cuh
@@ -108,7 +108,7 @@ _CCCL_HOST_DEVICE constexpr int sum(int head, Ts... tail)
 template <typename... Its>
 _CCCL_HOST_DEVICE constexpr auto loaded_bytes_per_iteration() -> int
 {
-  return (int{sizeof(value_t<Its>)} + ... + 0);
+  return (int{sizeof(it_value_t<Its>)} + ... + 0);
 }
 
 constexpr int bulk_copy_alignment     = 128;
@@ -144,7 +144,7 @@ struct policy_hub<RequiresStableAddress, ::cuda::std::tuple<RandomAccessIterator
   static constexpr bool all_contiguous =
     ::cuda::std::conjunction_v<THRUST_NS_QUALIFIER::is_contiguous_iterator<RandomAccessIteratorsIn>...>;
   static constexpr bool all_values_trivially_reloc =
-    ::cuda::std::conjunction_v<THRUST_NS_QUALIFIER::is_trivially_relocatable<value_t<RandomAccessIteratorsIn>>...>;
+    ::cuda::std::conjunction_v<THRUST_NS_QUALIFIER::is_trivially_relocatable<it_value_t<RandomAccessIteratorsIn>>...>;
 
   static constexpr bool can_memcpy = all_contiguous && all_values_trivially_reloc;
 
@@ -169,7 +169,7 @@ struct policy_hub<RequiresStableAddress, ::cuda::std::tuple<RandomAccessIterator
         async_policy::block_threads * async_policy::min_items_per_thread)
       > int{max_smem_per_block};
     static constexpr bool any_type_is_overalinged =
-      ((alignof(value_t<RandomAccessIteratorsIn>) > bulk_copy_alignment) || ...);
+      ((alignof(it_value_t<RandomAccessIteratorsIn>) > bulk_copy_alignment) || ...);
 
     static constexpr bool use_fallback =
       RequiresStableAddress || !can_memcpy || no_input_streams || exhaust_smem || any_type_is_overalinged;
diff --git a/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
index 57af475f1ba..093a17207e2 100644
--- a/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
@@ -849,6 +849,6 @@ struct policy_hub
 template <typename KeyInputIteratorT, typename ValueInputIteratorT = unsigned long long int*>
 using DeviceUniqueByKeyPolicy CCCL_DEPRECATED_BECAUSE("This class is considered an implementation detail and it will "
                                                       "be removed.") =
-  detail::unique_by_key::policy_hub<detail::value_t<KeyInputIteratorT>, detail::value_t<ValueInputIteratorT>>;
+  detail::unique_by_key::policy_hub<detail::it_value_t<KeyInputIteratorT>, detail::it_value_t<ValueInputIteratorT>>;
 
 CUB_NAMESPACE_END
diff --git a/cub/cub/iterator/arg_index_input_iterator.cuh b/cub/cub/iterator/arg_index_input_iterator.cuh
index c3916fcb2bb..be5a0b14086 100644
--- a/cub/cub/iterator/arg_index_input_iterator.cuh
+++ b/cub/cub/iterator/arg_index_input_iterator.cuh
@@ -46,7 +46,6 @@
 #include <cub/util_type.cuh>
 
 #include <thrust/iterator/iterator_facade.h>
-#include <thrust/iterator/iterator_traits.h>
 
 #if !_CCCL_COMPILER(NVRTC)
 #  include <ostream>
@@ -106,7 +105,7 @@ CUB_NAMESPACE_BEGIN
  */
 template <typename InputIteratorT,
           typename OffsetT      = ptrdiff_t,
-          typename OutputValueT = cub::detail::value_t<InputIteratorT>>
+          typename OutputValueT = detail::it_value_t<InputIteratorT>>
 class ArgIndexInputIterator
 {
 public:
diff --git a/cub/cub/iterator/cache_modified_input_iterator.cuh b/cub/cub/iterator/cache_modified_input_iterator.cuh
index ab97fae3525..e5d456aa011 100644
--- a/cub/cub/iterator/cache_modified_input_iterator.cuh
+++ b/cub/cub/iterator/cache_modified_input_iterator.cuh
@@ -47,7 +47,6 @@
 #  include <cuda/std/iterator>
 #else // ^^^ _CCCL_COMPILER(NVRTC) ^^^ // vvv !_CCCL_COMPILER(NVRTC) vvv
 #  include <thrust/iterator/iterator_facade.h>
-#  include <thrust/iterator/iterator_traits.h>
 
 #  include <iostream>
 #  include <iterator>
diff --git a/cub/cub/iterator/cache_modified_output_iterator.cuh b/cub/cub/iterator/cache_modified_output_iterator.cuh
index fa4e501b80f..e447c27a9c2 100644
--- a/cub/cub/iterator/cache_modified_output_iterator.cuh
+++ b/cub/cub/iterator/cache_modified_output_iterator.cuh
@@ -47,7 +47,6 @@
 #include <cub/thread/thread_store.cuh>
 
 #include <thrust/iterator/iterator_facade.h>
-#include <thrust/iterator/iterator_traits.h>
 
 #include <iosfwd>
 
diff --git a/cub/cub/iterator/tex_obj_input_iterator.cuh b/cub/cub/iterator/tex_obj_input_iterator.cuh
index b70819bde86..4be625c6fcc 100644
--- a/cub/cub/iterator/tex_obj_input_iterator.cuh
+++ b/cub/cub/iterator/tex_obj_input_iterator.cuh
@@ -48,7 +48,6 @@
 #include <cub/util_debug.cuh>
 
 #include <thrust/iterator/iterator_facade.h>
-#include <thrust/iterator/iterator_traits.h>
 
 #include <ostream>
 
diff --git a/cub/cub/thread/thread_load.cuh b/cub/cub/thread/thread_load.cuh
index 14b7269dfff..37c45d31536 100644
--- a/cub/cub/thread/thread_load.cuh
+++ b/cub/cub/thread/thread_load.cuh
@@ -106,7 +106,7 @@ enum CacheLoadModifier
  *   <b>[inferred]</b> The input's iterator type \iterator
  */
 template <CacheLoadModifier MODIFIER, typename RandomAccessIterator>
-_CCCL_DEVICE _CCCL_FORCEINLINE cub::detail::value_t<RandomAccessIterator> ThreadLoad(RandomAccessIterator itr);
+_CCCL_DEVICE _CCCL_FORCEINLINE cub::detail::it_value_t<RandomAccessIterator> ThreadLoad(RandomAccessIterator itr);
 
 //@}  end member group
 
@@ -306,7 +306,7 @@ _CUB_LOAD_ALL(LOAD_LDG, global.nc)
  * ThreadLoad definition for LOAD_DEFAULT modifier on iterator types
  */
 template <typename RandomAccessIterator>
-_CCCL_DEVICE _CCCL_FORCEINLINE cub::detail::value_t<RandomAccessIterator> ThreadLoad(
+_CCCL_DEVICE _CCCL_FORCEINLINE cub::detail::it_value_t<RandomAccessIterator> ThreadLoad(
   RandomAccessIterator itr, detail::constant_t<LOAD_DEFAULT> /*modifier*/, ::cuda::std::false_type /*is_pointer*/)
 {
   return *itr;
@@ -374,7 +374,7 @@ ThreadLoad(T const* ptr, detail::constant_t<MODIFIER> /*modifier*/, ::cuda::std:
 }
 
 template <CacheLoadModifier MODIFIER, typename RandomAccessIterator>
-_CCCL_DEVICE _CCCL_FORCEINLINE cub::detail::value_t<RandomAccessIterator> ThreadLoad(RandomAccessIterator itr)
+_CCCL_DEVICE _CCCL_FORCEINLINE cub::detail::it_value_t<RandomAccessIterator> ThreadLoad(RandomAccessIterator itr)
 {
   return ThreadLoad(
     itr, detail::constant_v<MODIFIER>, detail::bool_constant_v<::cuda::std::is_pointer_v<RandomAccessIterator>>);
diff --git a/cub/cub/thread/thread_search.cuh b/cub/cub/thread/thread_search.cuh
index 6471a55b351..dfd2e07f2c7 100644
--- a/cub/cub/thread/thread_search.cuh
+++ b/cub/cub/thread/thread_search.cuh
@@ -60,7 +60,7 @@ _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void MergePathSearch(
   OffsetT diagonal, AIteratorT a, BIteratorT b, OffsetT a_len, OffsetT b_len, CoordinateT& path_coordinate)
 {
   /// The value type of the input iterator
-  using T = cub::detail::value_t<AIteratorT>;
+  using T = cub::detail::it_value_t<AIteratorT>;
 
   OffsetT split_min = CUB_MAX(diagonal - b_len, 0);
   OffsetT split_max = CUB_MIN(diagonal, a_len);
diff --git a/cub/cub/util_type.cuh b/cub/cub/util_type.cuh
index 04d0062f048..073655e42dd 100644
--- a/cub/cub/util_type.cuh
+++ b/cub/cub/util_type.cuh
@@ -49,6 +49,7 @@
 #include <thrust/iterator/discard_iterator.h>
 
 #include <cuda/std/cstdint>
+#include <cuda/std/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
 
@@ -70,12 +71,6 @@ _CCCL_DIAG_PUSH
 _CCCL_DIAG_POP
 #endif // _CCCL_HAS_NVFP8()
 
-#if _CCCL_COMPILER(NVRTC)
-#  include <cuda/std/iterator>
-#else // ^^^ _CCCL_COMPILER(NVRTC) ^^^ // vvv !_CCCL_COMPILER(NVRTC) vvv
-#  include <iterator>
-#endif // _CCCL_COMPILER(NVRTC)
-
 CUB_NAMESPACE_BEGIN
 
 /******************************************************************************
@@ -85,16 +80,29 @@ CUB_NAMESPACE_BEGIN
 #ifndef _CCCL_DOXYGEN_INVOKED // Do not document
 namespace detail
 {
-//! Alias to the given iterator's value_type.
-// Aliases to std::iterator_traits, since users can specialize this template to provide traits for their iterators. We
-// only defer to the libcu++ implementation for NVRTC.
-template <typename Iterator>
-using value_t =
-#  if _CCCL_COMPILER(NVRTC)
-  typename ::cuda::std::iterator_traits<Iterator>::value_type;
-#  else // ^^^ _CCCL_COMPILER(NVRTC) ^^^ // vvv !_CCCL_COMPILER(NVRTC) vvv
-  typename std::iterator_traits<Iterator>::value_type;
-#  endif // !_CCCL_COMPILER(NVRTC)
+// the following iterator helpers are not named iter_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...>;
+};
 
 template <typename It, typename FallbackT, bool = ::cuda::std::is_void_v<::cuda::std::remove_pointer_t<It>>>
 struct non_void_value_impl
@@ -108,10 +116,10 @@ struct non_void_value_impl<It, FallbackT, false>
   // we consider thrust::discard_iterator's value_type as `void` as well, so users can switch from
   // cub::DiscardInputIterator to thrust::discard_iterator.
   using type =
-    ::cuda::std::_If<::cuda::std::is_void_v<value_t<It>>
-                       || ::cuda::std::is_same_v<value_t<It>, THRUST_NS_QUALIFIER::discard_iterator<>::value_type>,
+    ::cuda::std::_If<::cuda::std::is_void_v<it_value_t<It>>
+                       || ::cuda::std::is_same_v<it_value_t<It>, THRUST_NS_QUALIFIER::discard_iterator<>::value_type>,
                      FallbackT,
-                     value_t<It>>;
+                     it_value_t<It>>;
 };
 
 /**
diff --git a/cub/test/catch2_test_block_load.cu b/cub/test/catch2_test_block_load.cu
index da5db5bdb83..8aff2d8296c 100644
--- a/cub/test/catch2_test_block_load.cu
+++ b/cub/test/catch2_test_block_load.cu
@@ -49,7 +49,7 @@ template <int ItemsPerThread,
           typename OutputIteratorT>
 __global__ void kernel(cuda::std::true_type, InputIteratorT input, OutputIteratorT output, int num_items)
 {
-  using input_t      = cub::detail::value_t<InputIteratorT>;
+  using input_t      = cub::detail::it_value_t<InputIteratorT>;
   using block_load_t = cub::BlockLoad<input_t, ThreadsInBlock, ItemsPerThread, LoadAlgorithm>;
   using storage_t    = typename block_load_t::TempStorage;
 
diff --git a/cub/test/catch2_test_block_radix_sort.cuh b/cub/test/catch2_test_block_radix_sort.cuh
index 7cf6d03ddc2..4039f7d13bc 100644
--- a/cub/test/catch2_test_block_radix_sort.cuh
+++ b/cub/test/catch2_test_block_radix_sort.cuh
@@ -44,7 +44,7 @@ template <typename InputIteratorT,
 __global__ void
 kernel(ActionT action, InputIteratorT input, OutputIteratorT output, int begin_bit, int end_bit, bool striped)
 {
-  using key_t = cub::detail::value_t<InputIteratorT>;
+  using key_t = cub::detail::it_value_t<InputIteratorT>;
   using block_radix_sort_t =
     cub::BlockRadixSort<key_t, ThreadsInBlock, ItemsPerThread, cub::NullType, RadixBits, Memoize, Algorithm, ShmemConfig>;
 
@@ -121,8 +121,8 @@ __global__ void kernel(
   int end_bit,
   bool striped)
 {
-  using key_t   = cub::detail::value_t<InputKeyIteratorT>;
-  using value_t = cub::detail::value_t<InputValueIteratorT>;
+  using key_t   = cub::detail::it_value_t<InputKeyIteratorT>;
+  using value_t = cub::detail::it_value_t<InputValueIteratorT>;
   using block_radix_sort_t =
     cub::BlockRadixSort<key_t, ThreadsInBlock, ItemsPerThread, value_t, RadixBits, Memoize, Algorithm, ShmemConfig>;
 
diff --git a/cub/test/catch2_test_block_run_length_decode.cu b/cub/test/catch2_test_block_run_length_decode.cu
index cc90095f49f..f942e4f3987 100644
--- a/cub/test/catch2_test_block_run_length_decode.cu
+++ b/cub/test/catch2_test_block_run_length_decode.cu
@@ -74,8 +74,8 @@ public:
   static constexpr bool TEST_RELATIVE_OFFSETS = TEST_RELATIVE_OFFSETS_;
 
 private:
-  using RunItemT   = cub::detail::value_t<ItemItT>;
-  using RunLengthT = cub::detail::value_t<RunLengthsItT>;
+  using RunItemT   = cub::detail::it_value_t<ItemItT>;
+  using RunLengthT = cub::detail::it_value_t<RunLengthsItT>;
 
   using BlockRunOffsetScanT = cub::BlockScan<RunLengthT, BLOCK_DIM_X, cub::BLOCK_SCAN_RAKING, BLOCK_DIM_Y, BLOCK_DIM_Z>;
 
diff --git a/cub/test/catch2_test_block_store.cu b/cub/test/catch2_test_block_store.cu
index 447ec46ad2c..040f76a4f4a 100644
--- a/cub/test/catch2_test_block_store.cu
+++ b/cub/test/catch2_test_block_store.cu
@@ -57,7 +57,7 @@ template <typename InputIteratorT,
           cub::BlockStoreAlgorithm StoreAlgorithm>
 __global__ void kernel(std::integral_constant<bool, true>, InputIteratorT input, OutputIteratorT output, int num_items)
 {
-  using input_t       = cub::detail::value_t<InputIteratorT>;
+  using input_t       = cub::detail::it_value_t<InputIteratorT>;
   using block_store_t = cub::BlockStore<input_t, ThreadsInBlock, ItemsPerThread, StoreAlgorithm>;
   using storage_t     = typename block_store_t::TempStorage;
 
@@ -111,7 +111,7 @@ template <int ItemsPerThread,
           typename OutputIteratorT>
 void block_store(InputIteratorT input, OutputIteratorT output, int num_items)
 {
-  using input_t                       = cub::detail::value_t<InputIteratorT>;
+  using input_t                       = cub::detail::it_value_t<InputIteratorT>;
   using block_store_t                 = cub::BlockStore<input_t, ThreadsInBlock, ItemsPerThread, StoreAlgorithm>;
   using storage_t                     = typename block_store_t::TempStorage;
   constexpr bool sufficient_resources = sizeof(storage_t) <= cub::detail::max_smem_per_block;
diff --git a/cub/test/catch2_test_device_reduce.cu b/cub/test/catch2_test_device_reduce.cu
index da739a9dc35..1b9634b9ccb 100644
--- a/cub/test/catch2_test_device_reduce.cu
+++ b/cub/test/catch2_test_device_reduce.cu
@@ -149,7 +149,7 @@ C2H_TEST("Device reduce works with all device interfaces", "[reduce][device]", f
     // Run test
     c2h::device_vector<output_t> out_result(num_segments);
     auto d_out_it = thrust::raw_pointer_cast(out_result.data());
-    using init_t  = cub::detail::value_t<decltype(unwrap_it(d_out_it))>;
+    using init_t  = cub::detail::it_value_t<decltype(unwrap_it(d_out_it))>;
     device_reduce(unwrap_it(d_in_it), unwrap_it(d_out_it), num_items, reduction_op, init_t{});
 
     // Verify result
diff --git a/cub/test/catch2_test_device_reduce.cuh b/cub/test/catch2_test_device_reduce.cuh
index c43468a78b7..e1ed98a0b4b 100644
--- a/cub/test/catch2_test_device_reduce.cuh
+++ b/cub/test/catch2_test_device_reduce.cuh
@@ -272,7 +272,7 @@ inline void compute_host_reference(
     auto seg_end   = seg_begin + h_sizes_begin[segment];
     // TODO Should this be using cub accumulator t?
     h_data_out[segment] =
-      static_cast<cub::detail::value_t<ResultOutItT>>(std::accumulate(seg_begin, seg_end, init, reduction_op));
+      static_cast<cub::detail::it_value_t<ResultOutItT>>(std::accumulate(seg_begin, seg_end, init, reduction_op));
   }
 }
 
diff --git a/cub/test/catch2_test_device_scan.cu b/cub/test/catch2_test_device_scan.cu
index 4ee928079fe..32716ce667c 100644
--- a/cub/test/catch2_test_device_scan.cu
+++ b/cub/test/catch2_test_device_scan.cu
@@ -266,7 +266,7 @@ C2H_TEST("Device scan works with all device interfaces", "[scan][device]", full_
     // Run test
     c2h::device_vector<output_t> out_result(num_items);
     auto d_out_it = thrust::raw_pointer_cast(out_result.data());
-    using init_t  = cub::detail::value_t<decltype(unwrap_it(d_out_it))>;
+    using init_t  = cub::detail::it_value_t<decltype(unwrap_it(d_out_it))>;
     device_exclusive_scan(unwrap_it(d_in_it), unwrap_it(d_out_it), scan_op, init_t{}, num_items);
 
     // Verify result
@@ -301,7 +301,7 @@ C2H_TEST("Device scan works with all device interfaces", "[scan][device]", full_
     // Run test
     c2h::device_vector<output_t> out_result(num_items);
     auto d_out_it = thrust::raw_pointer_cast(out_result.data());
-    using init_t  = cub::detail::value_t<decltype(unwrap_it(d_out_it))>;
+    using init_t  = cub::detail::it_value_t<decltype(unwrap_it(d_out_it))>;
     c2h::device_vector<init_t> d_initial_value(1);
     d_initial_value[0]     = static_cast<init_t>(*unwrap_it(&init_value));
     auto future_init_value = cub::FutureValue<init_t>(thrust::raw_pointer_cast(d_initial_value.data()));
diff --git a/cub/test/catch2_test_device_scan.cuh b/cub/test/catch2_test_device_scan.cuh
index 62c731223da..36d027d4496 100644
--- a/cub/test/catch2_test_device_scan.cuh
+++ b/cub/test/catch2_test_device_scan.cuh
@@ -63,9 +63,9 @@ struct Mod2Equality
 template <typename InputIt, typename OutputIt, typename InitT, typename BinaryOp>
 void compute_exclusive_scan_reference(InputIt first, InputIt last, OutputIt result, InitT init, BinaryOp op)
 {
-  using value_t  = cub::detail::value_t<InputIt>;
+  using value_t  = cub::detail::it_value_t<InputIt>;
   using accum_t  = ::cuda::std::__accumulator_t<BinaryOp, value_t, InitT>;
-  using output_t = cub::detail::value_t<OutputIt>;
+  using output_t = cub::detail::it_value_t<OutputIt>;
   accum_t acc    = static_cast<accum_t>(init);
   for (; first != last; ++first)
   {
@@ -77,9 +77,9 @@ void compute_exclusive_scan_reference(InputIt first, InputIt last, OutputIt resu
 template <typename InputIt, typename OutputIt, typename BinaryOp, typename InitT>
 void compute_inclusive_scan_reference(InputIt first, InputIt last, OutputIt result, BinaryOp op, InitT init)
 {
-  using value_t  = cub::detail::value_t<InputIt>;
+  using value_t  = cub::detail::it_value_t<InputIt>;
   using accum_t  = ::cuda::std::__accumulator_t<BinaryOp, value_t, InitT>;
-  using output_t = cub::detail::value_t<OutputIt>;
+  using output_t = cub::detail::it_value_t<OutputIt>;
   accum_t acc    = static_cast<accum_t>(init);
   for (; first != last; ++first)
   {
@@ -103,9 +103,9 @@ void compute_exclusive_scan_by_key_reference(
   InitT init,
   std::size_t num_items)
 {
-  using value_t  = cub::detail::value_t<ValueInItT>;
+  using value_t  = cub::detail::it_value_t<ValueInItT>;
   using accum_t  = ::cuda::std::__accumulator_t<ScanOpT, value_t, InitT>;
-  using output_t = cub::detail::value_t<ValuesOutItT>;
+  using output_t = cub::detail::it_value_t<ValuesOutItT>;
 
   if (num_items > 0)
   {
@@ -154,9 +154,9 @@ void compute_inclusive_scan_by_key_reference(
   EqualityOpT equality_op,
   std::size_t num_items)
 {
-  using value_t  = cub::detail::value_t<ValueInItT>;
+  using value_t  = cub::detail::it_value_t<ValueInItT>;
   using accum_t  = ::cuda::std::__accumulator_t<ScanOpT, value_t, value_t>;
-  using output_t = cub::detail::value_t<ValuesOutItT>;
+  using output_t = cub::detail::it_value_t<ValuesOutItT>;
 
   for (std::size_t i = 0; i < num_items;)
   {
diff --git a/cub/test/catch2_test_device_scan_by_key.cu b/cub/test/catch2_test_device_scan_by_key.cu
index 5c92ae3300d..a1abe7cec52 100644
--- a/cub/test/catch2_test_device_scan_by_key.cu
+++ b/cub/test/catch2_test_device_scan_by_key.cu
@@ -232,7 +232,7 @@ C2H_TEST("Device scan works with all device interfaces", "[by_key][scan][device]
     // Run test
     c2h::device_vector<output_t> out_values(num_items);
     auto d_values_out_it = thrust::raw_pointer_cast(out_values.data());
-    using init_t         = cub::detail::value_t<decltype(unwrap_it(d_values_out_it))>;
+    using init_t         = cub::detail::it_value_t<decltype(unwrap_it(d_values_out_it))>;
     device_exclusive_scan_by_key(
       d_keys_it, unwrap_it(d_values_it), unwrap_it(d_values_out_it), scan_op, init_t{}, num_items, eq_op_t{});
 
diff --git a/cub/test/catch2_test_device_scan_iterators.cu b/cub/test/catch2_test_device_scan_iterators.cu
index 792b28a1854..a7e603b8741 100644
--- a/cub/test/catch2_test_device_scan_iterators.cu
+++ b/cub/test/catch2_test_device_scan_iterators.cu
@@ -170,7 +170,7 @@ C2H_TEST("Device scan works with iterators", "[scan][device]", iterator_type_lis
     // Run test
     c2h::device_vector<output_t> out_result(num_items);
     auto d_out_it = thrust::raw_pointer_cast(out_result.data());
-    using init_t  = cub::detail::value_t<decltype(unwrap_it(d_out_it))>;
+    using init_t  = cub::detail::it_value_t<decltype(unwrap_it(d_out_it))>;
     c2h::device_vector<init_t> d_initial_value(1);
     d_initial_value[0]     = static_cast<init_t>(init_value);
     auto future_init_value = cub::FutureValue<init_t>(thrust::raw_pointer_cast(d_initial_value.data()));
diff --git a/cub/test/catch2_test_device_segmented_reduce.cu b/cub/test/catch2_test_device_segmented_reduce.cu
index a3f0a73c8d2..064f09990a6 100644
--- a/cub/test/catch2_test_device_segmented_reduce.cu
+++ b/cub/test/catch2_test_device_segmented_reduce.cu
@@ -128,7 +128,7 @@ C2H_TEST("Device reduce works with all device interfaces", "[segmented][reduce][
     // Run test
     c2h::device_vector<output_t> out_result(num_segments);
     auto d_out_it = thrust::raw_pointer_cast(out_result.data());
-    using init_t  = cub::detail::value_t<decltype(unwrap_it(d_out_it))>;
+    using init_t  = cub::detail::it_value_t<decltype(unwrap_it(d_out_it))>;
     device_segmented_reduce(
       unwrap_it(d_in_it), unwrap_it(d_out_it), num_segments, d_offsets_it, d_offsets_it + 1, reduction_op, init_t{});
 
diff --git a/cub/test/catch2_test_device_select_common.cuh b/cub/test/catch2_test_device_select_common.cuh
index d41b0700512..cb6db23d0be 100644
--- a/cub/test/catch2_test_device_select_common.cuh
+++ b/cub/test/catch2_test_device_select_common.cuh
@@ -3,6 +3,8 @@
 
 #pragma once
 
+#include <cub/util_type.cuh>
+
 #include <thrust/iterator/constant_iterator.h>
 
 #include <cuda/std/type_traits>
@@ -59,7 +61,7 @@ struct modx_and_add_divy
 template <typename SelectedItT, typename RejectedItT>
 struct index_to_expected_partition_op
 {
-  using value_t = typename ::cuda::std::iterator_traits<SelectedItT>::value_type;
+  using value_t = cub::detail::it_value_t<SelectedItT>;
   SelectedItT expected_selected_it;
   RejectedItT expected_rejected_it;
   std::int64_t expected_num_selected;
diff --git a/cub/test/catch2_test_iterator.cu b/cub/test/catch2_test_iterator.cu
index 98b2855eb58..651e7d17315 100644
--- a/cub/test/catch2_test_iterator.cu
+++ b/cub/test/catch2_test_iterator.cu
@@ -36,6 +36,7 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH
 
 #include <cub/iterator/arg_index_input_iterator.cuh>
 #include <cub/iterator/cache_modified_input_iterator.cuh>
+#include <cub/iterator/cache_modified_output_iterator.cuh>
 #include <cub/iterator/tex_obj_input_iterator.cuh>
 #include <cub/util_allocator.cuh>
 #include <cub/util_type.cuh>
@@ -117,32 +118,35 @@ void test_iterator(InputIteratorT d_in, const c2h::host_vector<T>& h_reference)
   CHECK(d_in == h_itrs[1]);
 }
 
-using cache_modifiers =
-  c2h::enum_type_list<cub::CacheLoadModifier,
-                      cub::LOAD_DEFAULT,
-                      cub::LOAD_CA,
-                      cub::LOAD_CG,
-                      cub::LOAD_CS,
-                      cub::LOAD_CV,
-                      cub::LOAD_LDG,
-                      cub::LOAD_VOLATILE>;
-
-C2H_TEST("Test cache modified iterator", "[iterator]", types, cache_modifiers)
-{
-  using T                       = c2h::get<0, TestType>;
-  constexpr auto cache_modifier = c2h::get<1, TestType>::value;
-  constexpr int TEST_VALUES     = 11000;
-
-  c2h::device_vector<T> d_data(TEST_VALUES);
-  c2h::gen(C2H_SEED(1), d_data);
-  c2h::host_vector<T> h_data(d_data);
-
-  const auto h_reference = c2h::host_vector<T>{
-    h_data[0], h_data[100], h_data[1000], h_data[10000], h_data[1], h_data[21], h_data[11], h_data[0]};
-  test_iterator(
-    cub::CacheModifiedInputIterator<cache_modifier, T>(const_cast<const T*>(thrust::raw_pointer_cast(d_data.data()))),
-    h_reference);
-}
+static_assert(
+  ::cuda::std::is_void_v<cub::detail::it_value_t<cub::CacheModifiedOutputIterator<cub::STORE_DEFAULT, int>>>);
+
+// using cache_modifiers =
+//   c2h::enum_type_list<cub::CacheLoadModifier,
+//                       cub::LOAD_DEFAULT,
+//                       cub::LOAD_CA,
+//                       cub::LOAD_CG,
+//                       cub::LOAD_CS,
+//                       cub::LOAD_CV,
+//                       cub::LOAD_LDG,
+//                       cub::LOAD_VOLATILE>;
+//
+// C2H_TEST("Test cache modified iterator", "[iterator]", types, cache_modifiers)
+// {
+//   using T                       = c2h::get<0, TestType>;
+//   constexpr auto cache_modifier = c2h::get<1, TestType>::value;
+//   constexpr int TEST_VALUES     = 11000;
+//
+//   c2h::device_vector<T> d_data(TEST_VALUES);
+//   c2h::gen(C2H_SEED(1), d_data);
+//   c2h::host_vector<T> h_data(d_data);
+//
+//   const auto h_reference = c2h::host_vector<T>{
+//     h_data[0], h_data[100], h_data[1000], h_data[10000], h_data[1], h_data[21], h_data[11], h_data[0]};
+//   test_iterator(
+//     cub::CacheModifiedInputIterator<cache_modifier, T>(const_cast<const
+//     T*>(thrust::raw_pointer_cast(d_data.data()))), h_reference);
+// }
 
 template <typename T>
 struct transform_op_t
diff --git a/cub/test/catch2_test_vsmem.cu b/cub/test/catch2_test_vsmem.cu
index 53bfb00a28a..725ec2b594e 100644
--- a/cub/test/catch2_test_vsmem.cu
+++ b/cub/test/catch2_test_vsmem.cu
@@ -87,7 +87,7 @@ struct agent_dummy_algorithm_t
   static constexpr auto items_per_thread = ActivePolicyT::ITEMS_PER_THREAD;
   static constexpr auto tile_size        = block_threads * items_per_thread;
 
-  using item_t = cub::detail::value_t<InputIteratorT>;
+  using item_t = cub::detail::it_value_t<InputIteratorT>;
 
   using block_load_t = cub::BlockLoad<item_t, block_threads, items_per_thread, cub::BLOCK_LOAD_TRANSPOSE>;
 
@@ -193,7 +193,7 @@ void __global__ __launch_bounds__(
 template <typename InputIteratorT>
 struct device_dummy_algorithm_policy_t
 {
-  using item_t = cub::detail::value_t<InputIteratorT>;
+  using item_t = cub::detail::it_value_t<InputIteratorT>;
 
   static constexpr int FALLBACK_BLOCK_THREADS = 64;
 
@@ -219,7 +219,7 @@ template <typename InputIteratorT,
           typename PolicyHub = device_dummy_algorithm_policy_t<InputIteratorT>>
 struct dispatch_dummy_algorithm_t
 {
-  using item_t = cub::detail::value_t<InputIteratorT>;
+  using item_t = cub::detail::it_value_t<InputIteratorT>;
 
   /// Device-accessible allocation of temporary storage. When nullptr, the required
   /// allocation size is written to \p temp_storage_bytes and no work is done.
diff --git a/thrust/thrust/system/cuda/detail/scan.h b/thrust/thrust/system/cuda/detail/scan.h
index 3777422ff4a..5402437673f 100644
--- a/thrust/thrust/system/cuda/detail/scan.h
+++ b/thrust/thrust/system/cuda/detail/scan.h
@@ -120,7 +120,7 @@ _CCCL_HOST_DEVICE OutputIt inclusive_scan_n_impl(
   ScanOp scan_op)
 {
   using InputValueT = cub::detail::InputValue<InitValueT>;
-  using ValueT      = cub::detail::value_t<InputIt>;
+  using ValueT      = cub::detail::it_value_t<InputIt>;
   using AccumT      = ::cuda::std::__accumulator_t<ScanOp, ValueT, InitValueT>;
 
   using Dispatch32 =

From b77f3567e0cde66ed44590b4541a9cbf55aad461 Mon Sep 17 00:00:00 2001
From: Trent Nelson <trent@trent.me>
Date: Mon, 24 Feb 2025 12:09:53 -0800
Subject: [PATCH 05/17] Add C++ test for single-item-per-thread BlockScan Sum
 routines. (#3889)

* Use constexpr for mode selection.

Per Bernhard's feedback here:

    https://github.com/NVIDIA/cccl/pull/3829#discussion_r1961221710

* Implement C++ tests for single-item BlockScan *Sum routines.
---
 cub/test/catch2_test_block_scan.cu | 92 +++++++++++++++++++++++++++---
 1 file changed, 83 insertions(+), 9 deletions(-)

diff --git a/cub/test/catch2_test_block_scan.cu b/cub/test/catch2_test_block_scan.cu
index 32aa5d9f602..e81187330bb 100644
--- a/cub/test/catch2_test_block_scan.cu
+++ b/cub/test/catch2_test_block_scan.cu
@@ -68,6 +68,25 @@ __global__ void block_scan_kernel(T* in, T* out, ActionT action)
   }
 }
 
+template <cub::BlockScanAlgorithm Algorithm, int BlockDimX, int BlockDimY, int BlockDimZ, class T, class ActionT>
+__global__ void block_scan_single_kernel(T* in, T* out, ActionT action)
+{
+  using block_scan_t = cub::BlockScan<T, BlockDimX, Algorithm, BlockDimY, BlockDimZ>;
+  using storage_t    = typename block_scan_t::TempStorage;
+
+  __shared__ storage_t storage;
+
+  const int tid = static_cast<int>(cub::RowMajorTid(BlockDimX, BlockDimY, BlockDimZ));
+
+  T thread_data = in[tid];
+
+  block_scan_t scan(storage);
+
+  action(scan, thread_data);
+
+  out[tid] = thread_data;
+}
+
 template <cub::BlockScanAlgorithm Algorithm,
           int ItemsPerThread,
           int BlockDimX,
@@ -86,6 +105,18 @@ void block_scan(c2h::device_vector<T>& in, c2h::device_vector<T>& out, ActionT a
   REQUIRE(cudaSuccess == cudaDeviceSynchronize());
 }
 
+template <cub::BlockScanAlgorithm Algorithm, int BlockDimX, int BlockDimY, int BlockDimZ, class T, class ActionT>
+void block_scan_single(c2h::device_vector<T>& in, c2h::device_vector<T>& out, ActionT action)
+{
+  dim3 block_dims(BlockDimX, BlockDimY, BlockDimZ);
+
+  block_scan_single_kernel<Algorithm, BlockDimX, BlockDimY, BlockDimZ, T, ActionT>
+    <<<1, block_dims>>>(thrust::raw_pointer_cast(in.data()), thrust::raw_pointer_cast(out.data()), action);
+
+  REQUIRE(cudaSuccess == cudaPeekAtLastError());
+  REQUIRE(cudaSuccess == cudaDeviceSynchronize());
+}
+
 enum class scan_mode
 {
   exclusive,
@@ -98,7 +129,24 @@ struct sum_op_t
   template <int ItemsPerThread, class BlockScanT, class T>
   __device__ void operator()(BlockScanT& scan, T (&thread_data)[ItemsPerThread]) const
   {
-    if (Mode == scan_mode::exclusive)
+    if constexpr (Mode == scan_mode::exclusive)
+    {
+      scan.ExclusiveSum(thread_data, thread_data);
+    }
+    else
+    {
+      scan.InclusiveSum(thread_data, thread_data);
+    }
+  }
+};
+
+template <scan_mode Mode>
+struct sum_single_op_t
+{
+  template <class BlockScanT, class T>
+  __device__ void operator()(BlockScanT& scan, T& thread_data) const
+  {
+    if constexpr (Mode == scan_mode::exclusive)
     {
       scan.ExclusiveSum(thread_data, thread_data);
     }
@@ -133,7 +181,7 @@ struct min_op_t
   template <int ItemsPerThread, class BlockScanT>
   __device__ void operator()(BlockScanT& scan, int (&thread_data)[ItemsPerThread]) const
   {
-    if (Mode == scan_mode::exclusive)
+    if constexpr (Mode == scan_mode::exclusive)
     {
       scan.ExclusiveScan(thread_data, thread_data, ::cuda::minimum<>{});
     }
@@ -185,7 +233,7 @@ struct sum_aggregate_op_t
   {
     T block_aggregate{};
 
-    if (Mode == scan_mode::exclusive)
+    if constexpr (Mode == scan_mode::exclusive)
     {
       scan.ExclusiveSum(thread_data, thread_data, block_aggregate);
     }
@@ -232,7 +280,7 @@ struct sum_prefix_op_t
     const int tid = static_cast<int>(cub::RowMajorTid(blockDim.x, blockDim.y, blockDim.z));
     block_prefix_op_t prefix_op{tid, m_prefix};
 
-    if (Mode == scan_mode::exclusive)
+    if constexpr (Mode == scan_mode::exclusive)
     {
       scan.ExclusiveSum(thread_data, thread_data, prefix_op);
     }
@@ -273,7 +321,7 @@ struct min_prefix_op_t
     const int tid = static_cast<int>(cub::RowMajorTid(blockDim.x, blockDim.y, blockDim.z));
     block_prefix_op_t prefix_op{tid, m_prefix};
 
-    if (Mode == scan_mode::exclusive)
+    if constexpr (Mode == scan_mode::exclusive)
     {
       scan.ExclusiveScan(thread_data, thread_data, ::cuda::minimum<>{}, prefix_op);
     }
@@ -327,10 +375,11 @@ T host_scan(scan_mode mode, c2h::host_vector<T>& result, ScanOpT scan_op, T init
 
 using types = c2h::type_list<std::uint8_t, std::uint16_t, std::int32_t, std::int64_t>;
 // FIXME(bgruber): uchar3 fails the test, see #3835
-using vec_types        = c2h::type_list<ulonglong4, /*uchar3,*/ short2>;
-using block_dim_x      = c2h::enum_type_list<int, 17, 32, 65, 96>;
-using block_dim_yz     = c2h::enum_type_list<int, 1, 2>;
-using items_per_thread = c2h::enum_type_list<int, 1, 9>;
+using vec_types              = c2h::type_list<ulonglong4, /*uchar3,*/ short2>;
+using block_dim_x            = c2h::enum_type_list<int, 17, 32, 65, 96>;
+using block_dim_yz           = c2h::enum_type_list<int, 1, 2>;
+using items_per_thread       = c2h::enum_type_list<int, 1, 9>;
+using single_item_per_thread = c2h::enum_type_list<int, 1>;
 using algorithms =
   c2h::enum_type_list<cub::BlockScanAlgorithm,
                       cub::BlockScanAlgorithm::BLOCK_SCAN_RAKING,
@@ -378,6 +427,31 @@ C2H_TEST(
   REQUIRE_APPROX_EQ(h_out, d_out);
 }
 
+C2H_TEST("Block scan works with sum single",
+         "[scan][block]",
+         types,
+         block_dim_x,
+         block_dim_yz,
+         single_item_per_thread,
+         algorithm,
+         modes)
+{
+  using params = params_t<TestType>;
+  using type   = typename params::type;
+
+  c2h::device_vector<type> d_out(params::tile_size);
+  c2h::device_vector<type> d_in(params::tile_size);
+  c2h::gen(C2H_SEED(10), d_in);
+
+  block_scan_single<params::algorithm, params::block_dim_x, params::block_dim_y, params::block_dim_z>(
+    d_in, d_out, sum_single_op_t<params::mode>{});
+
+  c2h::host_vector<type> h_out = d_in;
+  host_scan(params::mode, h_out, std::plus<type>{});
+
+  REQUIRE_APPROX_EQ(h_out, d_out);
+}
+
 C2H_TEST("Block scan works with vec types", "[scan][block]", vec_types, algorithm, modes)
 {
   constexpr int items_per_thread              = 3;

From 048e236eab5592bd376b209933689c727250c332 Mon Sep 17 00:00:00 2001
From: Trent Nelson <trent@trent.me>
Date: Mon, 24 Feb 2025 12:10:54 -0800
Subject: [PATCH 06/17] Rename threads_in_block -> threads_per_block to be
 consistent with CUB. (#3919)

Closes #3911.
---
 .../experimental/block/_block_load_store.py   |   8 +-
 .../experimental/block/_block_merge_sort.py   |   6 +-
 .../experimental/block/_block_radix_sort.py   |  12 +-
 .../experimental/block/_block_reduce.py       |  12 +-
 .../experimental/block/_block_scan.py         |   8 +-
 .../cuda_cooperative/tests/test_block_load.py |  12 +-
 .../tests/test_block_merge_sort.py            |  30 ++--
 .../tests/test_block_radix_sort.py            |  38 +++---
 .../tests/test_block_reduce.py                | 128 ++++++++++--------
 .../tests/test_block_reduce_api.py            |  18 +--
 .../cuda_cooperative/tests/test_block_scan.py |  42 +++---
 .../tests/test_block_scan_api.py              |  16 +--
 .../tests/test_block_store.py                 |  12 +-
 13 files changed, 175 insertions(+), 167 deletions(-)

diff --git a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_load_store.py b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_load_store.py
index bc3f5d80ad1..200be5fa8d4 100644
--- a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_load_store.py
+++ b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_load_store.py
@@ -35,7 +35,7 @@
 }
 
 
-def load(dtype, threads_in_block, items_per_thread=1, algorithm="direct"):
+def load(dtype, threads_per_block, items_per_thread=1, algorithm="direct"):
     template = Algorithm(
         "BlockLoad",
         "Load",
@@ -58,7 +58,7 @@ def load(dtype, threads_in_block, items_per_thread=1, algorithm="direct"):
     specialization = template.specialize(
         {
             "T": dtype,
-            "BLOCK_DIM_X": threads_in_block,
+            "BLOCK_DIM_X": threads_per_block,
             "ITEMS_PER_THREAD": items_per_thread,
             "ALGORITHM": CUB_BLOCK_LOAD_ALGOS[algorithm],
         }
@@ -73,7 +73,7 @@ def load(dtype, threads_in_block, items_per_thread=1, algorithm="direct"):
     )
 
 
-def store(dtype, threads_in_block, items_per_thread=1, algorithm="direct"):
+def store(dtype, threads_per_block, items_per_thread=1, algorithm="direct"):
     template = Algorithm(
         "BlockStore",
         "Store",
@@ -96,7 +96,7 @@ def store(dtype, threads_in_block, items_per_thread=1, algorithm="direct"):
     specialization = template.specialize(
         {
             "T": dtype,
-            "BLOCK_DIM_X": threads_in_block,
+            "BLOCK_DIM_X": threads_per_block,
             "ITEMS_PER_THREAD": items_per_thread,
             "ALGORITHM": CUB_BLOCK_STORE_ALGOS[algorithm],
         }
diff --git a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_merge_sort.py b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_merge_sort.py
index 83911ffcac3..1659ef56517 100644
--- a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_merge_sort.py
+++ b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_merge_sort.py
@@ -19,7 +19,7 @@
 
 
 def merge_sort_keys(
-    dtype, threads_in_block, items_per_thread, compare_op, methods=None
+    dtype, threads_per_block, items_per_thread, compare_op, methods=None
 ):
     """Performs a block-wide merge sort over a :ref:`blocked arrangement <flexible-data-arrangement>` of keys.
 
@@ -49,7 +49,7 @@ def merge_sort_keys(
 
     Args:
         dtype: Numba data type of the keys to be sorted
-        threads_in_block: The number of threads in a block
+        threads_per_block: The number of threads in a block
         items_per_thread: The number of items each thread owns
         compare_op: Comparison function object which returns true if the first argument is ordered before the second one
 
@@ -82,7 +82,7 @@ def merge_sort_keys(
     specialization = template.specialize(
         {
             "KeyT": dtype,
-            "BLOCK_DIM_X": threads_in_block,
+            "BLOCK_DIM_X": threads_per_block,
             "ITEMS_PER_THREAD": items_per_thread,
             "Op": compare_op,
         }
diff --git a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_radix_sort.py b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_radix_sort.py
index ed18168393a..bd8e47ee877 100644
--- a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_radix_sort.py
+++ b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_radix_sort.py
@@ -16,7 +16,7 @@
 )
 
 
-def radix_sort_keys(dtype, threads_in_block, items_per_thread):
+def radix_sort_keys(dtype, threads_per_block, items_per_thread):
     """Performs an ascending block-wide radix sort over a :ref:`blocked arrangement <flexible-data-arrangement>` of keys.
 
     Example:
@@ -45,7 +45,7 @@ def radix_sort_keys(dtype, threads_in_block, items_per_thread):
 
     Args:
         dtype: Numba data type of the keys to be sorted
-        threads_in_block: The number of threads in a block
+        threads_per_block: The number of threads in a block
         items_per_thread: The number of items each thread owns
 
     Returns:
@@ -77,7 +77,7 @@ def radix_sort_keys(dtype, threads_in_block, items_per_thread):
     specialization = template.specialize(
         {
             "KeyT": dtype,
-            "BLOCK_DIM_X": threads_in_block,
+            "BLOCK_DIM_X": threads_per_block,
             "ITEMS_PER_THREAD": items_per_thread,
         }
     )
@@ -91,7 +91,7 @@ def radix_sort_keys(dtype, threads_in_block, items_per_thread):
     )
 
 
-def radix_sort_keys_descending(dtype, threads_in_block, items_per_thread):
+def radix_sort_keys_descending(dtype, threads_per_block, items_per_thread):
     """Performs an descending block-wide radix sort over a :ref:`blocked arrangement <flexible-data-arrangement>` of keys.
 
     Example:
@@ -120,7 +120,7 @@ def radix_sort_keys_descending(dtype, threads_in_block, items_per_thread):
 
     Args:
         dtype: Numba data type of the keys to be sorted
-        threads_in_block: The number of threads in a block
+        threads_per_block: The number of threads in a block
         items_per_thread: The number of items each thread owns
 
     Returns:
@@ -152,7 +152,7 @@ def radix_sort_keys_descending(dtype, threads_in_block, items_per_thread):
     specialization = template.specialize(
         {
             "KeyT": dtype,
-            "BLOCK_DIM_X": threads_in_block,
+            "BLOCK_DIM_X": threads_per_block,
             "ITEMS_PER_THREAD": items_per_thread,
         }
     )
diff --git a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_reduce.py b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_reduce.py
index 37e992fde4e..dbda854bf13 100644
--- a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_reduce.py
+++ b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_reduce.py
@@ -19,7 +19,7 @@
 )
 
 
-def reduce(dtype, threads_in_block, binary_op, items_per_thread=1, methods=None):
+def reduce(dtype, threads_per_block, binary_op, items_per_thread=1, methods=None):
     """Creates an operation that computes a block-wide reduction for thread\ :sub:`0` using the
     specified binary reduction functor.
 
@@ -35,7 +35,7 @@ def reduce(dtype, threads_in_block, binary_op, items_per_thread=1, methods=None)
 
     Args:
         dtype: Data type being reduced
-        threads_in_block: The number of threads in a block
+        threads_per_block: The number of threads in a block
         binary_op: Binary reduction function
         items_per_thread: The number of items each thread contributes to the reduction
         methods: A dict of methods for user-defined types
@@ -110,7 +110,7 @@ def reduce(dtype, threads_in_block, binary_op, items_per_thread=1, methods=None)
     specialization = template.specialize(
         {
             "T": dtype,
-            "BLOCK_DIM_X": threads_in_block,
+            "BLOCK_DIM_X": threads_per_block,
             "ITEMS_PER_THREAD": items_per_thread,
             "Op": binary_op,
         }
@@ -126,7 +126,7 @@ def reduce(dtype, threads_in_block, binary_op, items_per_thread=1, methods=None)
     )
 
 
-def sum(dtype, threads_in_block, items_per_thread=1, methods=None):
+def sum(dtype, threads_per_block, items_per_thread=1, methods=None):
     """Creates an operation that computes a block-wide reduction for thread\ :sub:`0` using
     addition (+) as the reduction operator.
 
@@ -142,7 +142,7 @@ def sum(dtype, threads_in_block, items_per_thread=1, methods=None):
 
     Args:
         dtype: Data type being reduced
-        threads_in_block: The number of threads in a block
+        threads_per_block: The number of threads in a block
         items_per_thread: The number of items each thread owns
         methods: A dict of methods for user-defined types
 
@@ -201,7 +201,7 @@ def sum(dtype, threads_in_block, items_per_thread=1, methods=None):
     specialization = template.specialize(
         {
             "T": dtype,
-            "BLOCK_DIM_X": threads_in_block,
+            "BLOCK_DIM_X": threads_per_block,
             "ITEMS_PER_THREAD": items_per_thread,
         }
     )
diff --git a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_scan.py b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_scan.py
index ff088657ef0..c1c098bb5ec 100644
--- a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_scan.py
+++ b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_scan.py
@@ -27,7 +27,7 @@
 
 def _scan(
     dtype: Type[numba.types.Number],
-    threads_in_block: int,
+    threads_per_block: int,
     items_per_thread: int = 1,
     mode: Literal["exclusive"] = "exclusive",
     scan_op: Literal["+"] = "+",
@@ -45,7 +45,7 @@ def _scan(
 
     specialization_kwds = {
         "T": dtype,
-        "BLOCK_DIM_X": threads_in_block,
+        "BLOCK_DIM_X": threads_per_block,
         "ALGORITHM": CUB_BLOCK_SCAN_ALGOS[algorithm],
     }
 
@@ -176,7 +176,7 @@ def _scan(
 
 def exclusive_sum(
     dtype: Type[numba.types.Number],
-    threads_in_block: int,
+    threads_per_block: int,
     items_per_thread: int = 1,
     prefix_op: Callable = None,
     algorithm: Literal["raking", "raking_memoize", "warp_scans"] = "raking",
@@ -186,7 +186,7 @@ def exclusive_sum(
     """
     return _scan(
         dtype=dtype,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
         mode="exclusive",
         scan_op="+",
diff --git a/python/cuda_cooperative/tests/test_block_load.py b/python/cuda_cooperative/tests/test_block_load.py
index fa2f311aab1..a151bb82411 100644
--- a/python/cuda_cooperative/tests/test_block_load.py
+++ b/python/cuda_cooperative/tests/test_block_load.py
@@ -15,7 +15,7 @@
 
 
 @pytest.mark.parametrize("T", [types.int8, types.int16, types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 256])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 256])
 @pytest.mark.parametrize("items_per_thread", [1, 3])
 @pytest.mark.parametrize(
     "algorithm",
@@ -28,15 +28,15 @@
         "warp_transpose_timesliced",
     ],
 )
-def test_block_load(T, threads_in_block, items_per_thread, algorithm):
-    block_load = cudax.block.load(T, threads_in_block, items_per_thread, algorithm)
+def test_block_load(T, threads_per_block, items_per_thread, algorithm):
+    block_load = cudax.block.load(T, threads_per_block, items_per_thread, algorithm)
     temp_storage_bytes = block_load.temp_storage_bytes
 
     if algorithm == "striped":
 
         @cuda.jit(device=True)
         def output_index(i):
-            return cuda.threadIdx.x + threads_in_block * i
+            return cuda.threadIdx.x + threads_per_block * i
     else:
 
         @cuda.jit(device=True)
@@ -52,11 +52,11 @@ def kernel(d_input, d_output):
             d_output[output_index(i)] = thread_data[i]
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
     h_input = random_int(items_per_tile, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
diff --git a/python/cuda_cooperative/tests/test_block_merge_sort.py b/python/cuda_cooperative/tests/test_block_merge_sort.py
index 8568bec1a49..432c79bab0e 100644
--- a/python/cuda_cooperative/tests/test_block_merge_sort.py
+++ b/python/cuda_cooperative/tests/test_block_merge_sort.py
@@ -16,15 +16,15 @@
 
 
 @pytest.mark.parametrize("T", [types.int8, types.int16, types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 256, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 256, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 3])
-def test_block_merge_sort(T, threads_in_block, items_per_thread):
+def test_block_merge_sort(T, threads_per_block, items_per_thread):
     def op(a, b):
         return a < b
 
     block_merge_sort = cudax.block.merge_sort_keys(
         dtype=T,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
         compare_op=op,
     )
@@ -42,11 +42,11 @@ def kernel(input, output):
             output[tid * items_per_thread + i] = thread_data[i]
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
     input = random_int(items_per_tile, dtype)
     d_input = cuda.to_device(input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
@@ -62,15 +62,15 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.int8, types.int16, types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 256, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 256, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 3])
-def test_block_merge_sort_descending(T, threads_in_block, items_per_thread):
+def test_block_merge_sort_descending(T, threads_per_block, items_per_thread):
     def op(a, b):
         return a > b
 
     block_merge_sort = cudax.block.merge_sort_keys(
         dtype=T,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
         compare_op=op,
     )
@@ -88,11 +88,11 @@ def kernel(input, output):
             output[tid * items_per_thread + i] = thread_data[i]
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
     input = random_int(items_per_tile, dtype)
     d_input = cuda.to_device(input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
@@ -109,15 +109,15 @@ def kernel(input, output):
 
 def test_block_merge_sort_user_defined_type():
     items_per_thread = 3
-    threads_in_block = 128
-    items_per_tile = threads_in_block * items_per_thread
+    threads_per_block = 128
+    items_per_tile = threads_per_block * items_per_thread
 
     def op(a, b):
         return a[0].real > b[0].real
 
     block_merge_sort = cudax.block.merge_sort_keys(
         dtype=numba.complex128,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
         compare_op=op,
     )
@@ -135,12 +135,12 @@ def kernel(input, output):
             output[tid * items_per_thread + i] = thread_data[i]
 
     dtype = np.complex128
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
     input = np.random.random(items_per_tile) + 1j * np.random.random(items_per_tile)
     input = input.astype(dtype)
     d_input = cuda.to_device(input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
diff --git a/python/cuda_cooperative/tests/test_block_radix_sort.py b/python/cuda_cooperative/tests/test_block_radix_sort.py
index 7e6e2642937..9b9b73ae851 100644
--- a/python/cuda_cooperative/tests/test_block_radix_sort.py
+++ b/python/cuda_cooperative/tests/test_block_radix_sort.py
@@ -15,13 +15,13 @@
 
 
 @pytest.mark.parametrize("T", [types.int8, types.int16, types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 256, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 256, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 3])
-def test_block_radix_sort_descending(T, threads_in_block, items_per_thread):
+def test_block_radix_sort_descending(T, threads_per_block, items_per_thread):
     begin_bit = numba.int32(0)
     end_bit = numba.int32(T.bitwidth)
     block_radix_sort = cudax.block.radix_sort_keys_descending(
-        dtype=T, threads_in_block=threads_in_block, items_per_thread=items_per_thread
+        dtype=T, threads_per_block=threads_per_block, items_per_thread=items_per_thread
     )
     temp_storage_bytes = block_radix_sort.temp_storage_bytes
 
@@ -37,11 +37,11 @@ def kernel(input, output):
             output[tid * items_per_thread + i] = thread_data[i]
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
     input = random_int(items_per_tile, dtype)
     d_input = cuda.to_device(input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
@@ -57,13 +57,13 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.int8, types.int16, types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 256, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 256, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 3])
-def test_block_radix_sort(T, threads_in_block, items_per_thread):
-    items_per_tile = threads_in_block * items_per_thread
+def test_block_radix_sort(T, threads_per_block, items_per_thread):
+    items_per_tile = threads_per_block * items_per_thread
 
     block_radix_sort = cudax.block.radix_sort_keys(
-        dtype=T, threads_in_block=threads_in_block, items_per_thread=items_per_thread
+        dtype=T, threads_per_block=threads_per_block, items_per_thread=items_per_thread
     )
     temp_storage_bytes = block_radix_sort.temp_storage_bytes
 
@@ -82,7 +82,7 @@ def kernel(input, output):
     input = random_int(items_per_tile, dtype)
     d_input = cuda.to_device(input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
@@ -99,12 +99,12 @@ def kernel(input, output):
 
 def test_block_radix_sort_overloads_work():
     T = numba.int32
-    threads_in_block = 128
+    threads_per_block = 128
     items_per_thread = 3
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
 
     block_radix_sort = cudax.block.radix_sort_keys(
-        dtype=T, threads_in_block=threads_in_block, items_per_thread=items_per_thread
+        dtype=T, threads_per_block=threads_per_block, items_per_thread=items_per_thread
     )
     temp_storage_bytes = block_radix_sort.temp_storage_bytes
 
@@ -123,7 +123,7 @@ def kernel(input, output):
     input = random_int(items_per_tile, dtype)
     d_input = cuda.to_device(input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
@@ -134,20 +134,20 @@ def kernel(input, output):
 
 def test_block_radix_sort_mangling():
     return  # TODO Return to linker issue
-    threads_in_block = 128
+    threads_per_block = 128
     items_per_thread = 3
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
 
     int_block_radix_sort = cudax.block.radix_sort_keys(
         dtype=numba.int32,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
     )
     int_temp_storage_bytes = int_block_radix_sort.temp_storage_bytes
 
     double_block_radix_sort = cudax.block.radix_sort_keys(
         dtype=numba.float64,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
     )
     double_temp_storage_bytes = double_block_radix_sort.temp_storage_bytes
@@ -182,7 +182,7 @@ def kernel(int_input, int_output, double_input, double_output):
     double_input = random_int(items_per_tile, "float64")
     d_double_input = cuda.to_device(double_input)
     d_double_output = cuda.device_array(items_per_tile, dtype="float64")
-    kernel[1, threads_in_block](
+    kernel[1, threads_per_block](
         d_int_input, d_int_output, d_double_input, d_double_output
     )
     cuda.synchronize()
diff --git a/python/cuda_cooperative/tests/test_block_reduce.py b/python/cuda_cooperative/tests/test_block_reduce.py
index d80c11e5b31..590298ff532 100644
--- a/python/cuda_cooperative/tests/test_block_reduce.py
+++ b/python/cuda_cooperative/tests/test_block_reduce.py
@@ -82,8 +82,8 @@ def impl_complex(context, builder, sig, args):
     return state._getvalue()
 
 
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
-def test_block_reduction_of_user_defined_type_without_temp_storage(threads_in_block):
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
+def test_block_reduction_of_user_defined_type_without_temp_storage(threads_per_block):
     def op(result_ptr, lhs_ptr, rhs_ptr):
         real_value = numba.int32(lhs_ptr[0].real + rhs_ptr[0].real)
         imag_value = numba.int32(lhs_ptr[0].imag + rhs_ptr[0].imag)
@@ -92,7 +92,7 @@ def op(result_ptr, lhs_ptr, rhs_ptr):
     block_reduce = cudax.block.reduce(
         dtype=complex_type,
         binary_op=op,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         methods={
             "construct": Complex.construct,
             "assign": Complex.assign,
@@ -102,20 +102,25 @@ def op(result_ptr, lhs_ptr, rhs_ptr):
     @cuda.jit(link=block_reduce.files)
     def kernel(input, output):
         block_output = block_reduce(
-            Complex(input[cuda.threadIdx.x], input[threads_in_block + cuda.threadIdx.x])
+            Complex(
+                input[cuda.threadIdx.x], input[threads_per_block + cuda.threadIdx.x]
+            )
         )
 
         if cuda.threadIdx.x == 0:
             output[0] = block_output.real
             output[1] = block_output.imag
 
-    h_input = random_int(2 * threads_in_block, "int32")
+    h_input = random_int(2 * threads_per_block, "int32")
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(2, dtype="int32")
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
-    h_expected = np.sum(h_input[:threads_in_block]), np.sum(h_input[threads_in_block:])
+    h_expected = (
+        np.sum(h_input[:threads_per_block]),
+        np.sum(h_input[threads_per_block:]),
+    )
 
     assert h_output[0] == h_expected[0]
     assert h_output[1] == h_expected[1]
@@ -127,8 +132,8 @@ def kernel(input, output):
     assert "STL" not in sass
 
 
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
-def test_block_reduction_of_user_defined_type(threads_in_block):
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
+def test_block_reduction_of_user_defined_type(threads_per_block):
     def op(result_ptr, lhs_ptr, rhs_ptr):
         real_value = numba.int32(lhs_ptr[0].real + rhs_ptr[0].real)
         imag_value = numba.int32(lhs_ptr[0].imag + rhs_ptr[0].imag)
@@ -137,7 +142,7 @@ def op(result_ptr, lhs_ptr, rhs_ptr):
     block_reduce = cudax.block.reduce(
         dtype=complex_type,
         binary_op=op,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         methods={
             "construct": Complex.construct,
             "assign": Complex.assign,
@@ -151,7 +156,7 @@ def kernel(input, output):
         block_output = block_reduce(
             temp_storage,
             Complex(
-                input[cuda.threadIdx.x], input[threads_in_block + cuda.threadIdx.x]
+                input[cuda.threadIdx.x], input[threads_per_block + cuda.threadIdx.x]
             ),
         )
 
@@ -159,13 +164,16 @@ def kernel(input, output):
             output[0] = block_output.real
             output[1] = block_output.imag
 
-    h_input = random_int(2 * threads_in_block, "int32")
+    h_input = random_int(2 * threads_per_block, "int32")
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(2, dtype="int32")
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
-    h_expected = np.sum(h_input[:threads_in_block]), np.sum(h_input[threads_in_block:])
+    h_expected = (
+        np.sum(h_input[:threads_per_block]),
+        np.sum(h_input[threads_per_block:]),
+    )
 
     assert h_output[0] == h_expected[0]
     assert h_output[1] == h_expected[1]
@@ -178,13 +186,13 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
-def test_block_reduction_of_integral_type(T, threads_in_block):
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
+def test_block_reduction_of_integral_type(T, threads_per_block):
     def op(a, b):
         return a if a < b else b
 
     block_reduce = cudax.block.reduce(
-        dtype=T, binary_op=op, threads_in_block=threads_in_block
+        dtype=T, binary_op=op, threads_per_block=threads_per_block
     )
     temp_storage_bytes = block_reduce.temp_storage_bytes
 
@@ -197,10 +205,10 @@ def kernel(input, output):
             output[0] = block_output
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    h_input = random_int(threads_in_block, dtype)
+    h_input = random_int(threads_per_block, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
     h_expected = np.min(h_input)
@@ -215,13 +223,13 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
-def test_block_reduction_valid(T, threads_in_block):
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
+def test_block_reduction_valid(T, threads_per_block):
     def op(a, b):
         return a if a < b else b
 
     block_reduce = cudax.block.reduce(
-        dtype=T, binary_op=op, threads_in_block=threads_in_block
+        dtype=T, binary_op=op, threads_per_block=threads_per_block
     )
     temp_storage_bytes = block_reduce.temp_storage_bytes
 
@@ -229,21 +237,21 @@ def op(a, b):
     def kernel(input, output):
         temp_storage = cuda.shared.array(shape=temp_storage_bytes, dtype="uint8")
         block_output = block_reduce(
-            temp_storage, input[cuda.threadIdx.x], threads_in_block / 2
+            temp_storage, input[cuda.threadIdx.x], threads_per_block / 2
         )
 
         if cuda.threadIdx.x == 0:
             output[0] = block_output
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    h_input = random_int(threads_in_block, dtype)
+    h_input = random_int(threads_per_block, dtype)
     h_input[-1] = 0
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
-    h_expected = np.min(h_input[: threads_in_block // 2])
+    h_expected = np.min(h_input[: threads_per_block // 2])
 
     assert h_output[0] == h_expected
 
@@ -255,16 +263,16 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 512, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 2, 4])
-def test_block_reduction_array_local(T, threads_in_block, items_per_thread):
+def test_block_reduction_array_local(T, threads_per_block, items_per_thread):
     def op(a, b):
         return a if a < b else b
 
     block_reduce = cudax.block.reduce(
         dtype=T,
         binary_op=op,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
     )
     temp_storage_bytes = block_reduce.temp_storage_bytes
@@ -275,7 +283,7 @@ def kernel(input, output):
         thread_items = cuda.local.array(shape=items_per_thread, dtype=T)
 
         for i in range(items_per_thread):
-            thread_items[i] = input[i * threads_in_block + cuda.threadIdx.x]
+            thread_items[i] = input[i * threads_per_block + cuda.threadIdx.x]
 
         block_output = block_reduce(temp_storage, thread_items)
 
@@ -283,10 +291,10 @@ def kernel(input, output):
             output[0] = block_output
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    h_input = random_int(items_per_thread * threads_in_block, dtype)
+    h_input = random_int(items_per_thread * threads_per_block, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
     h_expected = np.min(h_input)
@@ -301,16 +309,16 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 512, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 2, 4])
-def test_block_reduction_array_global(T, threads_in_block, items_per_thread):
+def test_block_reduction_array_global(T, threads_per_block, items_per_thread):
     def op(a, b):
         return a if a < b else b
 
     block_reduce = cudax.block.reduce(
         dtype=T,
         binary_op=op,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
     )
     temp_storage_bytes = block_reduce.temp_storage_bytes
@@ -331,10 +339,10 @@ def kernel(input, output):
             output[0] = block_output
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    h_input = random_int(items_per_thread * threads_in_block, dtype)
+    h_input = random_int(items_per_thread * threads_per_block, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
     h_expected = np.min(h_input)
@@ -349,9 +357,9 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
-def test_block_sum(T, threads_in_block):
-    block_reduce = cudax.block.sum(dtype=T, threads_in_block=threads_in_block)
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
+def test_block_sum(T, threads_per_block):
+    block_reduce = cudax.block.sum(dtype=T, threads_per_block=threads_per_block)
     temp_storage_bytes = block_reduce.temp_storage_bytes
 
     @cuda.jit(link=block_reduce.files)
@@ -363,10 +371,10 @@ def kernel(input, output):
             output[0] = block_output
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    h_input = random_int(threads_in_block, dtype)
+    h_input = random_int(threads_per_block, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
     h_expected = np.sum(h_input)
@@ -381,30 +389,30 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
-def test_block_sum_valid(T, threads_in_block):
-    block_reduce = cudax.block.sum(dtype=T, threads_in_block=threads_in_block)
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
+def test_block_sum_valid(T, threads_per_block):
+    block_reduce = cudax.block.sum(dtype=T, threads_per_block=threads_per_block)
     temp_storage_bytes = block_reduce.temp_storage_bytes
 
     @cuda.jit(link=block_reduce.files)
     def kernel(input, output):
         temp_storage = cuda.shared.array(shape=temp_storage_bytes, dtype="uint8")
         block_output = block_reduce(
-            temp_storage, input[cuda.threadIdx.x], numba.int32(threads_in_block / 2)
+            temp_storage, input[cuda.threadIdx.x], numba.int32(threads_per_block / 2)
         )
 
         if cuda.threadIdx.x == 0:
             output[0] = block_output
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    h_input = random_int(threads_in_block, dtype)
+    h_input = random_int(threads_per_block, dtype)
     h_input[-1] = 0
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
-    h_expected = np.sum(h_input[: threads_in_block // 2])
+    h_expected = np.sum(h_input[: threads_per_block // 2])
 
     assert h_output[0] == h_expected
 
@@ -416,11 +424,11 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 512, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 2, 4])
-def test_block_sum_array_local(T, threads_in_block, items_per_thread):
+def test_block_sum_array_local(T, threads_per_block, items_per_thread):
     block_reduce = cudax.block.sum(
-        dtype=T, threads_in_block=threads_in_block, items_per_thread=items_per_thread
+        dtype=T, threads_per_block=threads_per_block, items_per_thread=items_per_thread
     )
     temp_storage_bytes = block_reduce.temp_storage_bytes
 
@@ -430,7 +438,7 @@ def kernel(input, output):
         thread_items = cuda.local.array(shape=items_per_thread, dtype=T)
 
         for i in range(items_per_thread):
-            thread_items[i] = input[i * threads_in_block + cuda.threadIdx.x]
+            thread_items[i] = input[i * threads_per_block + cuda.threadIdx.x]
 
         block_output = block_reduce(temp_storage, thread_items)
 
@@ -438,10 +446,10 @@ def kernel(input, output):
             output[0] = block_output
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    h_input = random_int(items_per_thread * threads_in_block, dtype)
+    h_input = random_int(items_per_thread * threads_per_block, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
     h_expected = np.sum(h_input)
@@ -456,11 +464,11 @@ def kernel(input, output):
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 512, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 2, 4])
-def test_block_sum_array_global(T, threads_in_block, items_per_thread):
+def test_block_sum_array_global(T, threads_per_block, items_per_thread):
     block_reduce = cudax.block.sum(
-        dtype=T, threads_in_block=threads_in_block, items_per_thread=items_per_thread
+        dtype=T, threads_per_block=threads_per_block, items_per_thread=items_per_thread
     )
     temp_storage_bytes = block_reduce.temp_storage_bytes
 
@@ -480,10 +488,10 @@ def kernel(input, output):
             output[0] = block_output
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    h_input = random_int(items_per_thread * threads_in_block, dtype)
+    h_input = random_int(items_per_thread * threads_per_block, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
     h_expected = np.sum(h_input)
diff --git a/python/cuda_cooperative/tests/test_block_reduce_api.py b/python/cuda_cooperative/tests/test_block_reduce_api.py
index 8957f260598..85cb47ccef2 100644
--- a/python/cuda_cooperative/tests/test_block_reduce_api.py
+++ b/python/cuda_cooperative/tests/test_block_reduce_api.py
@@ -21,8 +21,8 @@ def test_block_reduction():
     def op(a, b):
         return a if a > b else b
 
-    threads_in_block = 128
-    block_reduce = cudax.block.reduce(numba.int32, threads_in_block, op)
+    threads_per_block = 128
+    block_reduce = cudax.block.reduce(numba.int32, threads_per_block, op)
 
     @cuda.jit(link=block_reduce.files)
     def kernel(input, output):
@@ -33,10 +33,10 @@ def kernel(input, output):
 
     # example-end reduce
 
-    h_input = np.random.randint(0, 42, threads_in_block, dtype=np.int32)
+    h_input = np.random.randint(0, 42, threads_per_block, dtype=np.int32)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=np.int32)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     h_output = d_output.copy_to_host()
     h_expected = np.max(h_input)
 
@@ -45,8 +45,8 @@ def kernel(input, output):
 
 def test_block_sum():
     # example-begin sum
-    threads_in_block = 128
-    block_sum = cudax.block.sum(numba.int32, threads_in_block)
+    threads_per_block = 128
+    block_sum = cudax.block.sum(numba.int32, threads_per_block)
 
     @cuda.jit(link=block_sum.files)
     def kernel(input, output):
@@ -57,10 +57,10 @@ def kernel(input, output):
 
     # example-end sum
 
-    h_input = np.ones(threads_in_block, dtype=np.int32)
+    h_input = np.ones(threads_per_block, dtype=np.int32)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(1, dtype=np.int32)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     h_output = d_output.copy_to_host()
 
-    assert h_output[0] == threads_in_block
+    assert h_output[0] == threads_per_block
diff --git a/python/cuda_cooperative/tests/test_block_scan.py b/python/cuda_cooperative/tests/test_block_scan.py
index 3f0c57ad95e..82ea4298481 100644
--- a/python/cuda_cooperative/tests/test_block_scan.py
+++ b/python/cuda_cooperative/tests/test_block_scan.py
@@ -27,19 +27,19 @@
 
 
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 2, 3, 4])
 @pytest.mark.parametrize("algorithm", ["raking", "raking_memoize", "warp_scans"])
-def test_block_exclusive_sum(T, threads_in_block, items_per_thread, algorithm):
-    if algorithm == "raking_memoize" and threads_in_block >= 512:
+def test_block_exclusive_sum(T, threads_per_block, items_per_thread, algorithm):
+    if algorithm == "raking_memoize" and threads_per_block >= 512:
         # We can hit CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES with raking_memoize in
-        # certain configurations, e.g.: 1024 threads_in_block, or 512
-        # threads_in_block, 3 items_per_thread, and T == uint64, etc.
-        pytest.skip("raking_memoize: skipping threads_in_block >= 512")
+        # certain configurations, e.g.: 1024 threads_per_block, or 512
+        # threads_per_block, 3 items_per_thread, and T == uint64, etc.
+        pytest.skip("raking_memoize: skipping threads_per_block >= 512")
 
     block_exclusive_sum = cudax.block.exclusive_sum(
         dtype=T,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
         algorithm=algorithm,
     )
@@ -60,11 +60,11 @@ def kernel(input, output):
             output[tid * items_per_thread + i] = thread_data[i]
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
     h_input = random_int(items_per_tile, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
@@ -130,17 +130,17 @@ def impl_block_prefix_callback_op(context, builder, sig, args):
     return state._getvalue()
 
 
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 2, 3, 4])
 @pytest.mark.parametrize("algorithm", ["raking", "raking_memoize", "warp_scans"])
-def test_block_exclusive_sum_prefix(threads_in_block, items_per_thread, algorithm):
-    if algorithm == "raking_memoize" and threads_in_block >= 512:
+def test_block_exclusive_sum_prefix(threads_per_block, items_per_thread, algorithm):
+    if algorithm == "raking_memoize" and threads_per_block >= 512:
         # We can hit CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES with raking_memoize in
-        # certain configurations, e.g.: 1024 threads_in_block, or 512
-        # threads_in_block, 3 items_per_thread, and T == uint64, etc.
-        pytest.skip("raking_memoize: skipping threads_in_block >= 512")
+        # certain configurations, e.g.: 1024 threads_per_block, or 512
+        # threads_per_block, 3 items_per_thread, and T == uint64, etc.
+        pytest.skip("raking_memoize: skipping threads_per_block >= 512")
 
-    tile_items = threads_in_block * items_per_thread
+    tile_items = threads_per_block * items_per_thread
     segment_size = 2 * 1024
     num_segments = 128
     num_elements = segment_size * num_segments
@@ -150,7 +150,7 @@ def test_block_exclusive_sum_prefix(threads_in_block, items_per_thread, algorith
     )
     block_exclusive_sum = cudax.block.exclusive_sum(
         dtype=numba.types.int32,
-        threads_in_block=threads_in_block,
+        threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
         prefix_op=prefix_op,
         algorithm=algorithm,
@@ -196,7 +196,7 @@ def kernel(input, output):
     h_input = np.arange(num_elements, dtype="int32")
     d_input = cuda.to_device(h_input)
     d_output = cuda.to_device(np.zeros(num_elements, dtype="int32"))
-    kernel[num_segments, threads_in_block](d_input, d_output)
+    kernel[num_segments, threads_per_block](d_input, d_output)
     cuda.synchronize()
     h_output = d_output.copy_to_host()
     h_reference = np.zeros(segment_size * num_segments, dtype="int32")
@@ -225,13 +225,13 @@ def kernel(input, output):
     assert "STL" not in sass
 
 
-@pytest.mark.parametrize("threads_in_block", [32, 64, 128, 256, 512, 1024])
+@pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [0, -1, -127])
 def test_block_scan_exclusive_sum_invalid_items_per_thread(
-    threads_in_block, items_per_thread
+    threads_per_block, items_per_thread
 ):
     with pytest.raises(ValueError):
-        cudax.block.exclusive_sum(numba.int32, threads_in_block, items_per_thread)
+        cudax.block.exclusive_sum(numba.int32, threads_per_block, items_per_thread)
 
 
 def test_block_scan_exclusive_sum_invalid_algorithm():
diff --git a/python/cuda_cooperative/tests/test_block_scan_api.py b/python/cuda_cooperative/tests/test_block_scan_api.py
index 3c2871bccb9..fa8976c1a60 100644
--- a/python/cuda_cooperative/tests/test_block_scan_api.py
+++ b/python/cuda_cooperative/tests/test_block_scan_api.py
@@ -19,11 +19,11 @@
 def test_block_exclusive_sum():
     # example-begin exclusive-sum
     items_per_thread = 4
-    threads_in_block = 128
+    threads_per_block = 128
 
     # Specialize exclusive sum for a 1D block of 128 threads owning 4 integer items each
     block_exclusive_sum = cudax.block.exclusive_sum(
-        numba.int32, threads_in_block, items_per_thread
+        numba.int32, threads_per_block, items_per_thread
     )
 
     # Link the exclusive sum to a CUDA kernel
@@ -43,11 +43,11 @@ def kernel(data):
 
     # example-end exclusive-sum
 
-    tile_size = threads_in_block * items_per_thread
+    tile_size = threads_per_block * items_per_thread
 
     h_keys = np.ones(tile_size, dtype=np.int32)
     d_keys = cuda.to_device(h_keys)
-    kernel[1, threads_in_block](d_keys)
+    kernel[1, threads_per_block](d_keys)
     h_keys = d_keys.copy_to_host()
     for i in range(tile_size):
         assert h_keys[i] == i
@@ -55,11 +55,11 @@ def kernel(data):
 
 def test_block_exclusive_sum_single_input_per_thread():
     # example-begin exclusive-sum-single-input-per-thread
-    threads_in_block = 128
+    threads_per_block = 128
 
     # Specialize exclusive sum for a 1D block of 128 threads.  Each thread
     # owns a single integer item.
-    block_exclusive_sum = cudax.block.exclusive_sum(numba.int32, threads_in_block)
+    block_exclusive_sum = cudax.block.exclusive_sum(numba.int32, threads_per_block)
 
     # Link the exclusive sum to a CUDA kernel
     @cuda.jit(link=block_exclusive_sum.files)
@@ -74,11 +74,11 @@ def kernel(data):
 
     # example-end exclusive-sum-single-input-per-thread
 
-    tile_size = threads_in_block
+    tile_size = threads_per_block
 
     h_keys = np.ones(tile_size, dtype=np.int32)
     d_keys = cuda.to_device(h_keys)
-    kernel[1, threads_in_block](d_keys)
+    kernel[1, threads_per_block](d_keys)
     h_keys = d_keys.copy_to_host()
     for i in range(tile_size):
         assert h_keys[i] == i
diff --git a/python/cuda_cooperative/tests/test_block_store.py b/python/cuda_cooperative/tests/test_block_store.py
index ce488f31a7b..973dcdd43c2 100644
--- a/python/cuda_cooperative/tests/test_block_store.py
+++ b/python/cuda_cooperative/tests/test_block_store.py
@@ -15,7 +15,7 @@
 
 
 @pytest.mark.parametrize("T", [types.int8, types.int16, types.uint32, types.uint64])
-@pytest.mark.parametrize("threads_in_block", [32, 128, 256])
+@pytest.mark.parametrize("threads_per_block", [32, 128, 256])
 @pytest.mark.parametrize("items_per_thread", [1, 3])
 @pytest.mark.parametrize(
     "algorithm",
@@ -28,15 +28,15 @@
         "warp_transpose_timesliced",
     ],
 )
-def test_block_store(T, threads_in_block, items_per_thread, algorithm):
-    block_store = cudax.block.store(T, threads_in_block, items_per_thread, algorithm)
+def test_block_store(T, threads_per_block, items_per_thread, algorithm):
+    block_store = cudax.block.store(T, threads_per_block, items_per_thread, algorithm)
     temp_storage_bytes = block_store.temp_storage_bytes
 
     if algorithm == "striped":
 
         @cuda.jit(device=True)
         def input_index(i):
-            return cuda.threadIdx.x + threads_in_block * i
+            return cuda.threadIdx.x + threads_per_block * i
     else:
 
         @cuda.jit(device=True)
@@ -52,11 +52,11 @@ def kernel(d_input, d_output):
         block_store(temp_storage, d_output, thread_data)
 
     dtype = NUMBA_TYPES_TO_NP[T]
-    items_per_tile = threads_in_block * items_per_thread
+    items_per_tile = threads_per_block * items_per_thread
     h_input = random_int(items_per_tile, dtype)
     d_input = cuda.to_device(h_input)
     d_output = cuda.device_array(items_per_tile, dtype=dtype)
-    kernel[1, threads_in_block](d_input, d_output)
+    kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()

From 39f183731e293bfd4e65afcd0390774829f6a798 Mon Sep 17 00:00:00 2001
From: Trent Nelson <trent@trent.me>
Date: Mon, 24 Feb 2025 16:24:37 -0800
Subject: [PATCH 07/17] Implement cuda.coopertive.block_scan.inclusive_sum().
 (#3921)

This new routine leverages the helper `block_scan._scan()` worker
function that was introduced in PR #3829.

Tests have been updated accordingly to test for both inclusive and
exclusive behavior, parameterized via the new `mode` argument.
---
 .../experimental/block/__init__.py            |   6 +-
 .../experimental/block/_block_scan.py         |  33 +++++-
 .../cuda_cooperative/tests/test_block_scan.py | 103 +++++++++++++-----
 3 files changed, 107 insertions(+), 35 deletions(-)

diff --git a/python/cuda_cooperative/cuda/cooperative/experimental/block/__init__.py b/python/cuda_cooperative/cuda/cooperative/experimental/block/__init__.py
index dd3161df731..3c006d1a461 100644
--- a/python/cuda_cooperative/cuda/cooperative/experimental/block/__init__.py
+++ b/python/cuda_cooperative/cuda/cooperative/experimental/block/__init__.py
@@ -9,13 +9,17 @@
     radix_sort_keys_descending,
 )
 from cuda.cooperative.experimental.block._block_reduce import reduce, sum
-from cuda.cooperative.experimental.block._block_scan import exclusive_sum
+from cuda.cooperative.experimental.block._block_scan import (
+    exclusive_sum,
+    inclusive_sum,
+)
 
 __all__ = [
     "merge_sort_keys",
     "reduce",
     "sum",
     "exclusive_sum",
+    "inclusive_sum",
     "radix_sort_keys",
     "radix_sort_keys_descending",
     "load",
diff --git a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_scan.py b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_scan.py
index c1c098bb5ec..e5847dd2ad5 100644
--- a/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_scan.py
+++ b/python/cuda_cooperative/cuda/cooperative/experimental/block/_block_scan.py
@@ -29,7 +29,7 @@ def _scan(
     dtype: Type[numba.types.Number],
     threads_per_block: int,
     items_per_thread: int = 1,
-    mode: Literal["exclusive"] = "exclusive",
+    mode: Literal["exclusive", "inclusive"] = "exclusive",
     scan_op: Literal["+"] = "+",
     block_prefix_callback_op: Callable = None,
     algorithm: Literal["raking", "raking_memoize", "warp_scans"] = "raking",
@@ -40,8 +40,12 @@ def _scan(
     if items_per_thread < 1:
         raise ValueError("items_per_thread must be greater than or equal to 1")
 
-    if mode != "exclusive":
-        raise ValueError(f"Unsupported mode: {mode}")
+    if mode == "exclusive":
+        cpp_func_prefix = "Exclusive"
+    else:
+        if mode != "inclusive":
+            raise ValueError(f"Unsupported mode: {mode}")
+        cpp_func_prefix = "Inclusive"
 
     specialization_kwds = {
         "T": dtype,
@@ -152,7 +156,7 @@ def _scan(
 
     template = Algorithm(
         "BlockScan",
-        "ExclusiveSum",
+        f"{cpp_func_prefix}Sum",
         "block_scan",
         ["cub/block/block_scan.cuh"],
         template_parameters,
@@ -193,3 +197,24 @@ def exclusive_sum(
         block_prefix_callback_op=prefix_op,
         algorithm=algorithm,
     )
+
+
+def inclusive_sum(
+    dtype: Type[numba.types.Number],
+    threads_per_block: int,
+    items_per_thread: int = 1,
+    prefix_op: Callable = None,
+    algorithm: Literal["raking", "raking_memoize", "warp_scans"] = "raking",
+) -> Callable:
+    """
+    Computes an inclusive block-wide prefix sum.
+    """
+    return _scan(
+        dtype=dtype,
+        threads_per_block=threads_per_block,
+        items_per_thread=items_per_thread,
+        mode="inclusive",
+        scan_op="+",
+        block_prefix_callback_op=prefix_op,
+        algorithm=algorithm,
+    )
diff --git a/python/cuda_cooperative/tests/test_block_scan.py b/python/cuda_cooperative/tests/test_block_scan.py
index 82ea4298481..c53d8f73565 100644
--- a/python/cuda_cooperative/tests/test_block_scan.py
+++ b/python/cuda_cooperative/tests/test_block_scan.py
@@ -29,46 +29,57 @@
 @pytest.mark.parametrize("T", [types.uint32, types.uint64])
 @pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 2, 3, 4])
+@pytest.mark.parametrize("mode", ["inclusive", "exclusive"])
 @pytest.mark.parametrize("algorithm", ["raking", "raking_memoize", "warp_scans"])
-def test_block_exclusive_sum(T, threads_per_block, items_per_thread, algorithm):
+def test_block_sum(T, threads_per_block, items_per_thread, mode, algorithm):
     if algorithm == "raking_memoize" and threads_per_block >= 512:
         # We can hit CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES with raking_memoize in
         # certain configurations, e.g.: 1024 threads_per_block, or 512
         # threads_per_block, 3 items_per_thread, and T == uint64, etc.
         pytest.skip("raking_memoize: skipping threads_per_block >= 512")
 
-    block_exclusive_sum = cudax.block.exclusive_sum(
+    if mode == "inclusive":
+        target_sum = cudax.block.inclusive_sum
+    else:
+        target_sum = cudax.block.exclusive_sum
+
+    block_sum = target_sum(
         dtype=T,
         threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
         algorithm=algorithm,
     )
-    temp_storage_bytes = block_exclusive_sum.temp_storage_bytes
+    temp_storage_bytes = block_sum.temp_storage_bytes
 
-    @cuda.jit(link=block_exclusive_sum.files)
+    @cuda.jit(link=block_sum.files)
     def kernel(input, output):
         tid = cuda.threadIdx.x
         temp_storage = cuda.shared.array(shape=temp_storage_bytes, dtype="uint8")
-        thread_data = cuda.local.array(shape=items_per_thread, dtype=dtype)
+        thread_data = cuda.local.array(shape=items_per_thread, dtype=T)
         for i in range(items_per_thread):
             thread_data[i] = input[tid * items_per_thread + i]
         if items_per_thread == 1:
-            thread_data[0] = block_exclusive_sum(temp_storage, thread_data[0])
+            thread_data[0] = block_sum(temp_storage, thread_data[0])
         else:
-            block_exclusive_sum(temp_storage, thread_data, thread_data)
+            block_sum(temp_storage, thread_data, thread_data)
         for i in range(items_per_thread):
             output[tid * items_per_thread + i] = thread_data[i]
 
-    dtype = NUMBA_TYPES_TO_NP[T]
+    dtype_np = NUMBA_TYPES_TO_NP[T]
     items_per_tile = threads_per_block * items_per_thread
-    h_input = random_int(items_per_tile, dtype)
+    h_input = random_int(items_per_tile, dtype_np)
     d_input = cuda.to_device(h_input)
-    d_output = cuda.device_array(items_per_tile, dtype=dtype)
+    d_output = cuda.device_array(items_per_tile, dtype=dtype_np)
     kernel[1, threads_per_block](d_input, d_output)
     cuda.synchronize()
 
     output = d_output.copy_to_host()
-    reference = np.cumsum(h_input) - h_input
+
+    if mode == "inclusive":
+        reference = np.cumsum(h_input)
+    else:
+        reference = np.cumsum(h_input) - h_input
+
     for i in range(items_per_tile):
         assert output[i] == reference[i]
 
@@ -132,8 +143,9 @@ def impl_block_prefix_callback_op(context, builder, sig, args):
 
 @pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [1, 2, 3, 4])
+@pytest.mark.parametrize("mode", ["inclusive", "exclusive"])
 @pytest.mark.parametrize("algorithm", ["raking", "raking_memoize", "warp_scans"])
-def test_block_exclusive_sum_prefix(threads_per_block, items_per_thread, algorithm):
+def test_block_sum_prefix(threads_per_block, items_per_thread, mode, algorithm):
     if algorithm == "raking_memoize" and threads_per_block >= 512:
         # We can hit CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES with raking_memoize in
         # certain configurations, e.g.: 1024 threads_per_block, or 512
@@ -148,16 +160,22 @@ def test_block_exclusive_sum_prefix(threads_per_block, items_per_thread, algorit
     prefix_op = cudax.StatefulFunction(
         BlockPrefixCallbackOp, block_prefix_callback_op_type
     )
-    block_exclusive_sum = cudax.block.exclusive_sum(
+
+    if mode == "inclusive":
+        target_sum = cudax.block.inclusive_sum
+    else:
+        target_sum = cudax.block.exclusive_sum
+
+    block_sum = target_sum(
         dtype=numba.types.int32,
         threads_per_block=threads_per_block,
         items_per_thread=items_per_thread,
         prefix_op=prefix_op,
         algorithm=algorithm,
     )
-    temp_storage_bytes = block_exclusive_sum.temp_storage_bytes
+    temp_storage_bytes = block_sum.temp_storage_bytes
 
-    @cuda.jit(link=block_exclusive_sum.files)
+    @cuda.jit(link=block_sum.files)
     def kernel(input, output):
         segment_offset = cuda.blockIdx.x * segment_size
         temp_storage = cuda.shared.array(shape=temp_storage_bytes, dtype="uint8")
@@ -178,13 +196,11 @@ def kernel(input, output):
                 )
 
             if items_per_thread == 1:
-                thread_output[0] = block_exclusive_sum(
+                thread_output[0] = block_sum(
                     temp_storage, thread_input[0], block_prefix_op
                 )
             else:
-                block_exclusive_sum(
-                    temp_storage, thread_input, thread_output, block_prefix_op
-                )
+                block_sum(temp_storage, thread_input, thread_output, block_prefix_op)
 
             for item in range(items_per_thread):
                 item_offset = tile_offset + cuda.threadIdx.x * items_per_thread + item
@@ -198,15 +214,30 @@ def kernel(input, output):
     d_output = cuda.to_device(np.zeros(num_elements, dtype="int32"))
     kernel[num_segments, threads_per_block](d_input, d_output)
     cuda.synchronize()
+
     h_output = d_output.copy_to_host()
     h_reference = np.zeros(segment_size * num_segments, dtype="int32")
-    for sid in range(num_segments):
-        h_reference[sid * segment_size] = 0
-        for i in range(1, segment_size):
-            h_reference[sid * segment_size + i] = (
-                h_reference[sid * segment_size + i - 1]
-                + h_input[sid * segment_size + i - 1]
-            )
+
+    if mode == "inclusive":
+        for sid in range(num_segments):
+            for i in range(segment_size):
+                if i == 0:
+                    h_reference[sid * segment_size + i] = h_input[
+                        sid * segment_size + i
+                    ]
+                else:
+                    h_reference[sid * segment_size + i] = (
+                        h_reference[sid * segment_size + i - 1]
+                        + h_input[sid * segment_size + i]
+                    )
+    else:
+        for sid in range(num_segments):
+            h_reference[sid * segment_size] = 0
+            for i in range(1, segment_size):
+                h_reference[sid * segment_size + i] = (
+                    h_reference[sid * segment_size + i - 1]
+                    + h_input[sid * segment_size + i - 1]
+                )
 
     for sid in range(num_segments):
         for i in range(segment_size):
@@ -227,13 +258,25 @@ def kernel(input, output):
 
 @pytest.mark.parametrize("threads_per_block", [32, 64, 128, 256, 512, 1024])
 @pytest.mark.parametrize("items_per_thread", [0, -1, -127])
+@pytest.mark.parametrize("mode", ["inclusive", "exclusive"])
 def test_block_scan_exclusive_sum_invalid_items_per_thread(
-    threads_per_block, items_per_thread
+    threads_per_block, items_per_thread, mode
 ):
+    if mode == "inclusive":
+        target_sum = cudax.block.inclusive_sum
+    else:
+        target_sum = cudax.block.exclusive_sum
+
     with pytest.raises(ValueError):
-        cudax.block.exclusive_sum(numba.int32, threads_per_block, items_per_thread)
+        target_sum(numba.int32, threads_per_block, items_per_thread)
+
 
+@pytest.mark.parametrize("mode", ["inclusive", "exclusive"])
+def test_block_scan_sum_invalid_algorithm(mode):
+    if mode == "inclusive":
+        target_sum = cudax.block.inclusive_sum
+    else:
+        target_sum = cudax.block.exclusive_sum
 
-def test_block_scan_exclusive_sum_invalid_algorithm():
     with pytest.raises(ValueError):
-        cudax.block.exclusive_sum(numba.int32, 128, algorithm="invalid_algorithm")
+        target_sum(numba.int32, 128, algorithm="invalid_algorithm")

From 1bb4e6afe6a165f5b1388fef076d820dc5af0db6 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Tue, 25 Feb 2025 19:07:54 +0100
Subject: [PATCH 08/17] Replace CUB macros in more places (#3930)

No SASS change on cub.test.device_radix_sort_pairs.lid_0 for SM86
---
 cub/cub/agent/agent_radix_sort_histogram.cuh  | 26 +++++++++----------
 cub/cub/agent/agent_radix_sort_upsweep.cuh    |  7 ++---
 cub/cub/block/block_radix_rank.cuh            |  7 ++---
 cub/cub/block/block_radix_sort.cuh            |  5 ++--
 .../device/dispatch/dispatch_radix_sort.cuh   | 25 +++++++++---------
 .../device/dispatch/kernels/radix_sort.cuh    | 10 ++++---
 .../example_device_partition_flagged.cu       | 14 +++++-----
 .../device/example_device_partition_if.cu     | 13 +++++-----
 .../device/example_device_select_flagged.cu   | 14 +++++-----
 .../device/example_device_select_if.cu        | 13 +++++-----
 .../device/example_device_select_unique.cu    | 14 +++++-----
 cub/test/catch2_radix_sort_helper.cuh         |  3 ++-
 cub/test/test_util.h                          |  8 +++---
 thrust/testing/tuple_reduce.cu                |  2 +-
 thrust/testing/tuple_scan.cu                  |  2 +-
 thrust/testing/tuple_sort.cu                  | 16 ++++++------
 thrust/testing/tuple_transform.cu             |  6 ++---
 thrust/testing/zip_function.cu                | 25 +++++++++---------
 18 files changed, 112 insertions(+), 98 deletions(-)

diff --git a/cub/cub/agent/agent_radix_sort_histogram.cuh b/cub/cub/agent/agent_radix_sort_histogram.cuh
index e2cc5ecac61..96825d581c0 100644
--- a/cub/cub/agent/agent_radix_sort_histogram.cuh
+++ b/cub/cub/agent/agent_radix_sort_histogram.cuh
@@ -51,6 +51,7 @@
 #include <cub/util_type.cuh>
 
 #include <cuda/ptx>
+#include <cuda/std/__algorithm_>
 
 CUB_NAMESPACE_BEGIN
 
@@ -66,7 +67,7 @@ struct AgentRadixSortHistogramPolicy
      * ID. However, lanes with the same ID in different warp use the same private
      * histogram. This arrangement helps reduce the degree of conflicts in atomic
      * operations. */
-    NUM_PARTS  = CUB_MAX(1, NOMINAL_4B_NUM_PARTS * 4 / CUB_MAX(sizeof(ComputeT), 4)),
+    NUM_PARTS  = _CUDA_VSTD::max(1, NOMINAL_4B_NUM_PARTS * 4 / _CUDA_VSTD::max(int{sizeof(ComputeT)}, 4)),
     RADIX_BITS = _RADIX_BITS,
   };
 };
@@ -94,16 +95,13 @@ template <typename AgentRadixSortHistogramPolicy,
 struct AgentRadixSortHistogram
 {
   // constants
-  enum
-  {
-    ITEMS_PER_THREAD = AgentRadixSortHistogramPolicy::ITEMS_PER_THREAD,
-    BLOCK_THREADS    = AgentRadixSortHistogramPolicy::BLOCK_THREADS,
-    TILE_ITEMS       = BLOCK_THREADS * ITEMS_PER_THREAD,
-    RADIX_BITS       = AgentRadixSortHistogramPolicy::RADIX_BITS,
-    RADIX_DIGITS     = 1 << RADIX_BITS,
-    MAX_NUM_PASSES   = (sizeof(KeyT) * 8 + RADIX_BITS - 1) / RADIX_BITS,
-    NUM_PARTS        = AgentRadixSortHistogramPolicy::NUM_PARTS,
-  };
+  static constexpr int ITEMS_PER_THREAD = AgentRadixSortHistogramPolicy::ITEMS_PER_THREAD;
+  static constexpr int BLOCK_THREADS    = AgentRadixSortHistogramPolicy::BLOCK_THREADS;
+  static constexpr int TILE_ITEMS       = BLOCK_THREADS * ITEMS_PER_THREAD;
+  static constexpr int RADIX_BITS       = AgentRadixSortHistogramPolicy::RADIX_BITS;
+  static constexpr int RADIX_DIGITS     = 1 << RADIX_BITS;
+  static constexpr int MAX_NUM_PASSES   = (sizeof(KeyT) * 8 + RADIX_BITS - 1) / RADIX_BITS;
+  static constexpr int NUM_PARTS        = AgentRadixSortHistogramPolicy::NUM_PARTS;
 
   using traits                 = radix::traits_t<KeyT>;
   using bit_ordered_type       = typename traits::bit_ordered_type;
@@ -210,7 +208,9 @@ struct AgentRadixSortHistogram
 #pragma unroll
     for (int current_bit = begin_bit, pass = 0; current_bit < end_bit; current_bit += RADIX_BITS, ++pass)
     {
-      int num_bits = CUB_MIN(RADIX_BITS, end_bit - current_bit);
+      // FIXME(bgruber): the following replacement changes SASS for cub.test.device_radix_sort_pairs.lid_0
+      // const int num_bits = _CUDA_VSTD::min(+RADIX_BITS, end_bit - current_bit);
+      const int num_bits = CUB_MIN(+RADIX_BITS, end_bit - current_bit);
 #pragma unroll
       for (int u = 0; u < ITEMS_PER_THREAD; ++u)
       {
@@ -258,7 +258,7 @@ struct AgentRadixSortHistogram
 
       // Process the tiles.
       OffsetT portion_offset = portion * MAX_PORTION_SIZE;
-      OffsetT portion_size   = CUB_MIN(MAX_PORTION_SIZE, num_items - portion_offset);
+      OffsetT portion_size   = _CUDA_VSTD::min(MAX_PORTION_SIZE, num_items - portion_offset);
       for (OffsetT offset = blockIdx.x * TILE_ITEMS; offset < portion_size; offset += TILE_ITEMS * gridDim.x)
       {
         OffsetT tile_offset = portion_offset + offset;
diff --git a/cub/cub/agent/agent_radix_sort_upsweep.cuh b/cub/cub/agent/agent_radix_sort_upsweep.cuh
index c8fa2ed841e..da34a0b4e75 100644
--- a/cub/cub/agent/agent_radix_sort_upsweep.cuh
+++ b/cub/cub/agent/agent_radix_sort_upsweep.cuh
@@ -53,6 +53,7 @@
 #include <cub/warp/warp_reduce.cuh>
 
 #include <cuda/ptx>
+#include <cuda/std/__algorithm_>
 
 CUB_NAMESPACE_BEGIN
 
@@ -160,17 +161,17 @@ struct AgentRadixSortUpsweep
     PACKING_RATIO     = sizeof(PackedCounter) / sizeof(DigitCounter),
     LOG_PACKING_RATIO = Log2<PACKING_RATIO>::VALUE,
 
-    LOG_COUNTER_LANES = CUB_MAX(0, int(RADIX_BITS) - int(LOG_PACKING_RATIO)),
+    LOG_COUNTER_LANES = _CUDA_VSTD::max(0, int(RADIX_BITS) - int(LOG_PACKING_RATIO)),
     COUNTER_LANES     = 1 << LOG_COUNTER_LANES,
 
     // To prevent counter overflow, we must periodically unpack and aggregate the
     // digit counters back into registers.  Each counter lane is assigned to a
     // warp for aggregation.
 
-    LANES_PER_WARP = CUB_MAX(1, (COUNTER_LANES + WARPS - 1) / WARPS),
+    LANES_PER_WARP = _CUDA_VSTD::max(1, (COUNTER_LANES + WARPS - 1) / WARPS),
 
     // Unroll tiles in batches without risk of counter overflow
-    UNROLL_COUNT      = CUB_MIN(64, 255 / KEYS_PER_THREAD),
+    UNROLL_COUNT      = _CUDA_VSTD::min(64, 255 / KEYS_PER_THREAD),
     UNROLLED_ELEMENTS = UNROLL_COUNT * TILE_ITEMS,
   };
 
diff --git a/cub/cub/block/block_radix_rank.cuh b/cub/cub/block/block_radix_rank.cuh
index 425ba1c576b..199c87a8380 100644
--- a/cub/cub/block/block_radix_rank.cuh
+++ b/cub/cub/block/block_radix_rank.cuh
@@ -49,6 +49,7 @@
 #include <cub/util_type.cuh>
 
 #include <cuda/ptx>
+#include <cuda/std/__algorithm_>
 #include <cuda/std/cstdint>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
@@ -242,7 +243,7 @@ private:
     LOG_PACKING_RATIO = Log2<PACKING_RATIO>::VALUE,
 
     // Always at least one lane
-    LOG_COUNTER_LANES = CUB_MAX((int(RADIX_BITS) - int(LOG_PACKING_RATIO)), 0),
+    LOG_COUNTER_LANES = _CUDA_VSTD::max(RADIX_BITS - LOG_PACKING_RATIO, 0),
     COUNTER_LANES     = 1 << LOG_COUNTER_LANES,
 
     // The number of packed counters per thread (plus one for padding)
@@ -254,7 +255,7 @@ public:
   enum
   {
     /// Number of bin-starting offsets tracked per thread
-    BINS_TRACKED_PER_THREAD = CUB_MAX(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS),
+    BINS_TRACKED_PER_THREAD = _CUDA_VSTD::max(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS),
   };
 
 private:
@@ -587,7 +588,7 @@ public:
   enum
   {
     /// Number of bin-starting offsets tracked per thread
-    BINS_TRACKED_PER_THREAD = CUB_MAX(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS),
+    BINS_TRACKED_PER_THREAD = _CUDA_VSTD::max(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS),
   };
 
 private:
diff --git a/cub/cub/block/block_radix_sort.cuh b/cub/cub/block/block_radix_sort.cuh
index 397930de643..3e18fe3372a 100644
--- a/cub/cub/block/block_radix_sort.cuh
+++ b/cub/cub/block/block_radix_sort.cuh
@@ -50,6 +50,7 @@
 #include <cub/util_ptx.cuh>
 #include <cub/util_type.cuh>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/type_traits>
 
 CUB_NAMESPACE_BEGIN
@@ -431,7 +432,7 @@ private:
     // Radix sorting passes
     while (true)
     {
-      int pass_bits = CUB_MIN(RADIX_BITS, end_bit - begin_bit);
+      int pass_bits = _CUDA_VSTD::min(RADIX_BITS, end_bit - begin_bit);
       auto digit_extractor =
         traits::template digit_extractor<fundamental_digit_extractor_t>(begin_bit, pass_bits, decomposer);
 
@@ -510,7 +511,7 @@ public:
     // Radix sorting passes
     while (true)
     {
-      int pass_bits = CUB_MIN(RADIX_BITS, end_bit - begin_bit);
+      int pass_bits = _CUDA_VSTD::min(RADIX_BITS, end_bit - begin_bit);
       auto digit_extractor =
         traits::template digit_extractor<fundamental_digit_extractor_t>(begin_bit, pass_bits, decomposer);
 
diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh
index 880a2cacce3..f967c21cf58 100644
--- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh
+++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh
@@ -53,6 +53,7 @@
 
 #include <thrust/system/cuda/detail/core/triple_chevron_launch.h>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/type_traits>
 
 #include <iterator>
@@ -275,7 +276,7 @@ struct DispatchRadixSort
     cudaError error = cudaSuccess;
     do
     {
-      int pass_bits = CUB_MIN(pass_config.radix_bits, (end_bit - current_bit));
+      int pass_bits = _CUDA_VSTD::min(pass_config.radix_bits, end_bit - current_bit);
 
 // Log upsweep_kernel configuration
 #ifdef CUB_DEBUG_LOG
@@ -447,7 +448,7 @@ struct DispatchRadixSort
         max_downsweep_grid_size = (downsweep_config.sm_occupancy * sm_count) * CUB_SUBSCRIPTION_FACTOR(0);
 
         even_share.DispatchInit(
-          num_items, max_downsweep_grid_size, CUB_MAX(downsweep_config.tile_size, upsweep_config.tile_size));
+          num_items, max_downsweep_grid_size, _CUDA_VSTD::max(downsweep_config.tile_size, upsweep_config.tile_size));
 
       } while (0);
       return error;
@@ -472,8 +473,8 @@ struct DispatchRadixSort
     constexpr PortionOffsetT PORTION_SIZE = ((1 << 28) - 1) / ONESWEEP_TILE_ITEMS * ONESWEEP_TILE_ITEMS;
     int num_passes                        = ::cuda::ceil_div(end_bit - begin_bit, RADIX_BITS);
     OffsetT num_portions                  = static_cast<OffsetT>(::cuda::ceil_div(num_items, PORTION_SIZE));
-    PortionOffsetT max_num_blocks =
-      ::cuda::ceil_div(static_cast<int>(CUB_MIN(num_items, static_cast<OffsetT>(PORTION_SIZE))), ONESWEEP_TILE_ITEMS);
+    PortionOffsetT max_num_blocks         = ::cuda::ceil_div(
+      static_cast<int>(_CUDA_VSTD::min(num_items, static_cast<OffsetT>(PORTION_SIZE))), ONESWEEP_TILE_ITEMS);
 
     size_t value_size         = KEYS_ONLY ? 0 : sizeof(ValueT);
     size_t allocation_sizes[] = {
@@ -611,11 +612,11 @@ struct DispatchRadixSort
 
       for (int current_bit = begin_bit, pass = 0; current_bit < end_bit; current_bit += RADIX_BITS, ++pass)
       {
-        int num_bits = CUB_MIN(end_bit - current_bit, RADIX_BITS);
+        int num_bits = _CUDA_VSTD::min(end_bit - current_bit, RADIX_BITS);
         for (OffsetT portion = 0; portion < num_portions; ++portion)
         {
-          PortionOffsetT portion_num_items = static_cast<PortionOffsetT>(
-            CUB_MIN(num_items - portion * PORTION_SIZE, static_cast<OffsetT>(PORTION_SIZE)));
+          PortionOffsetT portion_num_items =
+            static_cast<PortionOffsetT>(_CUDA_VSTD::min(num_items - portion * PORTION_SIZE, OffsetT{PORTION_SIZE}));
 
           PortionOffsetT num_blocks = ::cuda::ceil_div(portion_num_items, ONESWEEP_TILE_ITEMS);
 
@@ -777,7 +778,7 @@ struct DispatchRadixSort
       }
 
       // Get maximum spine length
-      int max_grid_size = CUB_MAX(pass_config.max_downsweep_grid_size, alt_pass_config.max_downsweep_grid_size);
+      int max_grid_size = _CUDA_VSTD::max(pass_config.max_downsweep_grid_size, alt_pass_config.max_downsweep_grid_size);
       int spine_length  = (max_grid_size * pass_config.radix_digits) + pass_config.scan_config.tile_size;
 
       // Temporary storage allocation requirements
@@ -812,7 +813,7 @@ struct DispatchRadixSort
       int num_passes         = ::cuda::ceil_div(num_bits, pass_config.radix_bits);
       bool is_num_passes_odd = num_passes & 1;
       int max_alt_passes     = (num_passes * pass_config.radix_bits) - num_bits;
-      int alt_end_bit        = CUB_MIN(end_bit, begin_bit + (max_alt_passes * alt_pass_config.radix_bits));
+      int alt_end_bit        = _CUDA_VSTD::min(end_bit, begin_bit + (max_alt_passes * alt_pass_config.radix_bits));
 
       // Alias the temporary storage allocations
       OffsetT* d_spine = static_cast<OffsetT*>(allocations[0]);
@@ -1241,7 +1242,7 @@ struct DispatchSegmentedRadixSort
     cudaError error = cudaSuccess;
     do
     {
-      int pass_bits = CUB_MIN(pass_config.radix_bits, (end_bit - current_bit));
+      int pass_bits = _CUDA_VSTD::min(pass_config.radix_bits, (end_bit - current_bit));
 
 // Log kernel configuration
 #ifdef CUB_DEBUG_LOG
@@ -1381,10 +1382,10 @@ struct DispatchSegmentedRadixSort
       int radix_bits         = ActivePolicyT::SegmentedPolicy::RADIX_BITS;
       int alt_radix_bits     = ActivePolicyT::AltSegmentedPolicy::RADIX_BITS;
       int num_bits           = end_bit - begin_bit;
-      int num_passes         = CUB_MAX(::cuda::ceil_div(num_bits, radix_bits), 1);
+      int num_passes         = _CUDA_VSTD::max(::cuda::ceil_div(num_bits, radix_bits), 1); // num_bits may be zero
       bool is_num_passes_odd = num_passes & 1;
       int max_alt_passes     = (num_passes * radix_bits) - num_bits;
-      int alt_end_bit        = CUB_MIN(end_bit, begin_bit + (max_alt_passes * alt_radix_bits));
+      int alt_end_bit        = _CUDA_VSTD::min(end_bit, begin_bit + (max_alt_passes * alt_radix_bits));
 
       DoubleBuffer<KeyT> d_keys_remaining_passes(
         (is_overwrite_okay || is_num_passes_odd) ? d_keys.Alternate() : static_cast<KeyT*>(allocations[0]),
diff --git a/cub/cub/device/dispatch/kernels/radix_sort.cuh b/cub/cub/device/dispatch/kernels/radix_sort.cuh
index 65327df5ec4..d49cc11b2c8 100644
--- a/cub/cub/device/dispatch/kernels/radix_sort.cuh
+++ b/cub/cub/device/dispatch/kernels/radix_sort.cuh
@@ -22,6 +22,8 @@
 #include <cub/device/dispatch/dispatch_common.cuh>
 #include <cub/grid/grid_even_share.cuh>
 
+#include <cuda/std/__algorithm_>
+
 CUB_NAMESPACE_BEGIN
 
 /******************************************************************************
@@ -98,8 +100,8 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? int(ChainedPolicyT::ActivePolicy::AltUp
 
   enum
   {
-    TILE_ITEMS = CUB_MAX(ActiveUpsweepPolicyT::BLOCK_THREADS * ActiveUpsweepPolicyT::ITEMS_PER_THREAD,
-                         ActiveDownsweepPolicyT::BLOCK_THREADS * ActiveDownsweepPolicyT::ITEMS_PER_THREAD)
+    TILE_ITEMS = _CUDA_VSTD::max(ActiveUpsweepPolicyT::BLOCK_THREADS * ActiveUpsweepPolicyT::ITEMS_PER_THREAD,
+                                 ActiveDownsweepPolicyT::BLOCK_THREADS * ActiveDownsweepPolicyT::ITEMS_PER_THREAD)
   };
 
   // Parameterize AgentRadixSortUpsweep type for the current configuration
@@ -258,8 +260,8 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? int(ChainedPolicyT::ActivePolicy::AltDo
 
   enum
   {
-    TILE_ITEMS = CUB_MAX(ActiveUpsweepPolicyT::BLOCK_THREADS * ActiveUpsweepPolicyT::ITEMS_PER_THREAD,
-                         ActiveDownsweepPolicyT::BLOCK_THREADS * ActiveDownsweepPolicyT::ITEMS_PER_THREAD)
+    TILE_ITEMS = _CUDA_VSTD::max(ActiveUpsweepPolicyT::BLOCK_THREADS * ActiveUpsweepPolicyT::ITEMS_PER_THREAD,
+                                 ActiveDownsweepPolicyT::BLOCK_THREADS * ActiveDownsweepPolicyT::ITEMS_PER_THREAD)
   };
 
   // Parameterize AgentRadixSortDownsweep type for the current configuration
diff --git a/cub/examples/device/example_device_partition_flagged.cu b/cub/examples/device/example_device_partition_flagged.cu
index 2b498d94d53..bb7840672c1 100644
--- a/cub/examples/device/example_device_partition_flagged.cu
+++ b/cub/examples/device/example_device_partition_flagged.cu
@@ -43,6 +43,8 @@
 #include <cub/device/device_partition.cuh>
 #include <cub/util_allocator.cuh>
 
+#include <cuda/std/limits>
+
 #include "../../test/test_util.h"
 #include <stdio.h>
 
@@ -65,20 +67,18 @@ CachingDeviceAllocator g_allocator(true); // Caching allocator for device memory
  */
 void Initialize(int* h_in, unsigned char* h_flags, int num_items, int max_segment)
 {
-  unsigned short max_short = (unsigned short) -1;
-
   int key = 0;
   int i   = 0;
   while (i < num_items)
   {
     // Select number of repeating occurrences
-    unsigned short repeat;
-    RandomBits(repeat);
-    repeat = (unsigned short) ((float(repeat) * (float(max_segment) / float(max_short))));
-    repeat = CUB_MAX(1, repeat);
+    unsigned short bits;
+    RandomBits(bits);
+    const int repeat = cuda::std::max(
+      1, static_cast<int>(bits * (static_cast<float>(max_segment) / cuda::std::numeric_limits<unsigned short>::max())));
 
     int j = i;
-    while (j < CUB_MIN(i + repeat, num_items))
+    while (j < cuda::std::min(i + repeat, num_items))
     {
       h_flags[j] = 0;
       h_in[j]    = key;
diff --git a/cub/examples/device/example_device_partition_if.cu b/cub/examples/device/example_device_partition_if.cu
index f7c2ea86d54..3daa0c54ed5 100644
--- a/cub/examples/device/example_device_partition_if.cu
+++ b/cub/examples/device/example_device_partition_if.cu
@@ -43,6 +43,8 @@
 #include <cub/device/device_partition.cuh>
 #include <cub/util_allocator.cuh>
 
+#include <cuda/std/limits>
+
 #include "../../test/test_util.h"
 #include <stdio.h>
 
@@ -84,14 +86,13 @@ void Initialize(int* h_in, int num_items, int max_segment)
   while (i < num_items)
   {
     // Randomly select number of repeating occurrences uniformly from [1..max_segment]
-    unsigned short max_short = (unsigned short) -1;
-    unsigned short repeat;
-    RandomBits(repeat);
-    repeat = (unsigned short) ((float(repeat) * (float(max_segment) / float(max_short))));
-    repeat = CUB_MAX(1, repeat);
+    unsigned short bits;
+    RandomBits(bits);
+    const int repeat = cuda::std::max(
+      1, static_cast<int>(bits * (static_cast<float>(max_segment) / cuda::std::numeric_limits<unsigned short>::max())));
 
     int j = i;
-    while (j < CUB_MIN(i + repeat, num_items))
+    while (j < cuda::std::min(i + repeat, num_items))
     {
       h_in[j] = key;
       j++;
diff --git a/cub/examples/device/example_device_select_flagged.cu b/cub/examples/device/example_device_select_flagged.cu
index 7f0ad094ce3..f50bee6a667 100644
--- a/cub/examples/device/example_device_select_flagged.cu
+++ b/cub/examples/device/example_device_select_flagged.cu
@@ -43,6 +43,8 @@
 #include <cub/device/device_select.cuh>
 #include <cub/util_allocator.cuh>
 
+#include <cuda/std/limits>
+
 #include "../../test/test_util.h"
 #include <stdio.h>
 
@@ -65,20 +67,18 @@ CachingDeviceAllocator g_allocator(true); // Caching allocator for device memory
  */
 void Initialize(int* h_in, unsigned char* h_flags, int num_items, int max_segment)
 {
-  unsigned short max_short = (unsigned short) -1;
-
   int key = 0;
   int i   = 0;
   while (i < num_items)
   {
     // Select number of repeating occurrences
-    unsigned short repeat;
-    RandomBits(repeat);
-    repeat = (unsigned short) ((float(repeat) * (float(max_segment) / float(max_short))));
-    repeat = CUB_MAX(1, repeat);
+    unsigned short bits;
+    RandomBits(bits);
+    const int repeat = cuda::std::max(
+      1, static_cast<int>(bits * (static_cast<float>(max_segment) / cuda::std::numeric_limits<unsigned short>::max())));
 
     int j = i;
-    while (j < CUB_MIN(i + repeat, num_items))
+    while (j < cuda::std::min(i + repeat, num_items))
     {
       h_flags[j] = 0;
       h_in[j]    = key;
diff --git a/cub/examples/device/example_device_select_if.cu b/cub/examples/device/example_device_select_if.cu
index ce070347d72..3da212ddf2b 100644
--- a/cub/examples/device/example_device_select_if.cu
+++ b/cub/examples/device/example_device_select_if.cu
@@ -43,6 +43,8 @@
 #include <cub/device/device_select.cuh>
 #include <cub/util_allocator.cuh>
 
+#include <cuda/std/limits>
+
 #include "../../test/test_util.h"
 #include <stdio.h>
 
@@ -84,14 +86,13 @@ void Initialize(int* h_in, int num_items, int max_segment)
   while (i < num_items)
   {
     // Randomly select number of repeating occurrences uniformly from [1..max_segment]
-    unsigned short max_short = (unsigned short) -1;
-    unsigned short repeat;
-    RandomBits(repeat);
-    repeat = (unsigned short) ((float(repeat) * (float(max_segment) / float(max_short))));
-    repeat = CUB_MAX(1, repeat);
+    unsigned short bits;
+    RandomBits(bits);
+    const int repeat = cuda::std::max(
+      1, static_cast<int>(bits * (static_cast<float>(max_segment) / cuda::std::numeric_limits<unsigned short>::max())));
 
     int j = i;
-    while (j < CUB_MIN(i + repeat, num_items))
+    while (j < cuda::std::min(i + repeat, num_items))
     {
       h_in[j] = key;
       j++;
diff --git a/cub/examples/device/example_device_select_unique.cu b/cub/examples/device/example_device_select_unique.cu
index dbd6a2462f1..1c361683d65 100644
--- a/cub/examples/device/example_device_select_unique.cu
+++ b/cub/examples/device/example_device_select_unique.cu
@@ -43,6 +43,8 @@
 #include <cub/device/device_select.cuh>
 #include <cub/util_allocator.cuh>
 
+#include <cuda/std/limits>
+
 #include "../../test/test_util.h"
 #include <stdio.h>
 
@@ -69,14 +71,14 @@ void Initialize(int* h_in, int num_items, int max_segment)
   while (i < num_items)
   {
     // Randomly select number of repeating occurrences uniformly from [1..max_segment]
-    unsigned short max_short = (unsigned short) -1;
-    unsigned short repeat;
-    RandomBits(repeat);
-    repeat = (unsigned short) ((float(repeat) * (float(max_segment) / float(max_short))));
-    repeat = CUB_MAX(1, repeat);
+    unsigned short bits;
+    RandomBits(bits);
+    const int repeat = cuda::std::max(
+      1, static_cast<int>(bits * (static_cast<float>(max_segment) / cuda::std::numeric_limits<unsigned short>::max())));
+    ;
 
     int j = i;
-    while (j < CUB_MIN(i + repeat, num_items))
+    while (j < cuda::std::min(i + repeat, num_items))
     {
       h_in[j] = key;
       j++;
diff --git a/cub/test/catch2_radix_sort_helper.cuh b/cub/test/catch2_radix_sort_helper.cuh
index 821927a0c05..89e525c1372 100644
--- a/cub/test/catch2_radix_sort_helper.cuh
+++ b/cub/test/catch2_radix_sort_helper.cuh
@@ -40,6 +40,7 @@
 #include <thrust/scan.h>
 #include <thrust/sequence.h>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/bit>
 #include <cuda/std/functional>
 #include <cuda/type_traits>
@@ -477,7 +478,7 @@ struct radix_offset_scan_op_t
   __host__ __device__ OffsetT operator()(OffsetT a, OffsetT b) const
   {
     const OffsetT sum = a + b;
-    return CUB_MIN(sum, num_items);
+    return _CUDA_VSTD::min(sum, num_items);
   }
 };
 
diff --git a/cub/test/test_util.h b/cub/test/test_util.h
index 88552be32ac..ee1ef03f140 100644
--- a/cub/test/test_util.h
+++ b/cub/test/test_util.h
@@ -45,6 +45,8 @@
 
 #include <thrust/iterator/discard_iterator.h>
 
+#include <cuda/std/__algorithm_>
+
 #include <cfloat>
 #include <cmath>
 #include <cstddef>
@@ -505,8 +507,8 @@ void RandomBits(K& key, int entropy_reduction = 0, int begin_bit = 0, int end_bi
       int current_bit = j * WORD_BYTES * 8;
 
       unsigned int word = 0xffffffff;
-      word &= 0xffffffff << CUB_MAX(0, begin_bit - current_bit);
-      word &= 0xffffffff >> CUB_MAX(0, (current_bit + (WORD_BYTES * 8)) - end_bit);
+      word &= 0xffffffff << ::cuda::std::max(0, begin_bit - current_bit);
+      word &= 0xffffffff >> ::cuda::std::max(0, (current_bit + (WORD_BYTES * 8)) - end_bit);
 
       for (int i = 0; i <= entropy_reduction; i++)
       {
@@ -1392,7 +1394,7 @@ void InitializeSegments(OffsetT num_items, int num_segments, OffsetT* h_segment_
 
     OffsetT segment_length = RandomValue((expected_segment_length * 2) + 1);
     offset += segment_length;
-    offset = CUB_MIN(offset, num_items);
+    offset = ::cuda::std::min(offset, num_items);
   }
   h_segment_offsets[num_segments] = num_items;
 
diff --git a/thrust/testing/tuple_reduce.cu b/thrust/testing/tuple_reduce.cu
index c7f7a42f7dc..2bcfab751a7 100644
--- a/thrust/testing/tuple_reduce.cu
+++ b/thrust/testing/tuple_reduce.cu
@@ -38,7 +38,7 @@ struct TestTupleReduce
 
     // zip up the data
     host_vector<tuple<T, T>> h_tuples(n);
-    transform(h_t1.begin(), h_t1.end(), h_t2.begin(), h_tuples.begin(), MakeTupleFunctor());
+    thrust::transform(h_t1.begin(), h_t1.end(), h_t2.begin(), h_tuples.begin(), MakeTupleFunctor());
 
     // copy to device
     device_vector<tuple<T, T>> d_tuples = h_tuples;
diff --git a/thrust/testing/tuple_scan.cu b/thrust/testing/tuple_scan.cu
index 298c2d87389..2a22e452734 100644
--- a/thrust/testing/tuple_scan.cu
+++ b/thrust/testing/tuple_scan.cu
@@ -42,7 +42,7 @@ struct TestTupleScan
 
     // initialize input
     host_vector<tuple<T, T>> h_input(n);
-    transform(h_t1.begin(), h_t1.end(), h_t2.begin(), h_input.begin(), MakeTupleFunctor());
+    thrust::transform(h_t1.begin(), h_t1.end(), h_t2.begin(), h_input.begin(), MakeTupleFunctor());
     device_vector<tuple<T, T>> d_input = h_input;
 
     // allocate output
diff --git a/thrust/testing/tuple_sort.cu b/thrust/testing/tuple_sort.cu
index 0263ebbd8d5..14156235f0d 100644
--- a/thrust/testing/tuple_sort.cu
+++ b/thrust/testing/tuple_sort.cu
@@ -37,30 +37,30 @@ struct TestTupleStableSort
 
     // zip up the data
     host_vector<tuple<T, T>> h_tuples(n);
-    transform(h_keys.begin(), h_keys.end(), h_values.begin(), h_tuples.begin(), MakeTupleFunctor());
+    thrust::transform(h_keys.begin(), h_keys.end(), h_values.begin(), h_tuples.begin(), MakeTupleFunctor());
 
     // copy to device
     device_vector<tuple<T, T>> d_tuples = h_tuples;
 
     // sort on host
-    stable_sort(h_tuples.begin(), h_tuples.end());
+    thrust::stable_sort(h_tuples.begin(), h_tuples.end());
 
     // sort on device
-    stable_sort(d_tuples.begin(), d_tuples.end());
+    thrust::stable_sort(d_tuples.begin(), d_tuples.end());
 
-    ASSERT_EQUAL(true, is_sorted(d_tuples.begin(), d_tuples.end()));
+    ASSERT_EQUAL(true, thrust::is_sorted(d_tuples.begin(), d_tuples.end()));
 
     // select keys
-    transform(h_tuples.begin(), h_tuples.end(), h_keys.begin(), GetFunctor<0>());
+    thrust::transform(h_tuples.begin(), h_tuples.end(), h_keys.begin(), GetFunctor<0>());
 
     device_vector<T> d_keys(h_keys.size());
-    transform(d_tuples.begin(), d_tuples.end(), d_keys.begin(), GetFunctor<0>());
+    thrust::transform(d_tuples.begin(), d_tuples.end(), d_keys.begin(), GetFunctor<0>());
 
     // select values
-    transform(h_tuples.begin(), h_tuples.end(), h_values.begin(), GetFunctor<1>());
+    thrust::transform(h_tuples.begin(), h_tuples.end(), h_values.begin(), GetFunctor<1>());
 
     device_vector<T> d_values(h_values.size());
-    transform(d_tuples.begin(), d_tuples.end(), d_values.begin(), GetFunctor<1>());
+    thrust::transform(d_tuples.begin(), d_tuples.end(), d_values.begin(), GetFunctor<1>());
 
     ASSERT_ALMOST_EQUAL(h_keys, d_keys);
     ASSERT_ALMOST_EQUAL(h_values, d_values);
diff --git a/thrust/testing/tuple_transform.cu b/thrust/testing/tuple_transform.cu
index 2b6dcc96c3f..a77bf2ad647 100644
--- a/thrust/testing/tuple_transform.cu
+++ b/thrust/testing/tuple_transform.cu
@@ -36,7 +36,7 @@ struct TestTupleTransform
 
     // zip up the data
     host_vector<tuple<T, T>> h_tuples(n);
-    transform(h_t1.begin(), h_t1.end(), h_t2.begin(), h_tuples.begin(), MakeTupleFunctor());
+    thrust::transform(h_t1.begin(), h_t1.end(), h_t2.begin(), h_tuples.begin(), MakeTupleFunctor());
 
     // copy to device
     device_vector<tuple<T, T>> d_tuples = h_tuples;
@@ -44,10 +44,10 @@ struct TestTupleTransform
     device_vector<T> d_t1(n), d_t2(n);
 
     // select 0th
-    transform(d_tuples.begin(), d_tuples.end(), d_t1.begin(), GetFunctor<0>());
+    thrust::transform(d_tuples.begin(), d_tuples.end(), d_t1.begin(), GetFunctor<0>());
 
     // select 1st
-    transform(d_tuples.begin(), d_tuples.end(), d_t2.begin(), GetFunctor<1>());
+    thrust::transform(d_tuples.begin(), d_tuples.end(), d_t2.begin(), GetFunctor<1>());
 
     ASSERT_ALMOST_EQUAL(h_t1, d_t1);
     ASSERT_ALMOST_EQUAL(h_t2, d_t2);
diff --git a/thrust/testing/zip_function.cu b/thrust/testing/zip_function.cu
index 279960592b5..2b2276b9c93 100644
--- a/thrust/testing/zip_function.cu
+++ b/thrust/testing/zip_function.cu
@@ -60,19 +60,20 @@ struct TestZipFunctionTransform
     device_vector<T> d_result_zip(n);
 
     // Tuple base case
-    transform(make_zip_iterator(h_data0.begin(), h_data1.begin(), h_data2.begin()),
-              make_zip_iterator(h_data0.end(), h_data1.end(), h_data2.end()),
-              h_result_tuple.begin(),
-              SumThreeTuple{});
+
+    thrust::transform(make_zip_iterator(h_data0.begin(), h_data1.begin(), h_data2.begin()),
+                      make_zip_iterator(h_data0.end(), h_data1.end(), h_data2.end()),
+                      h_result_tuple.begin(),
+                      SumThreeTuple{});
     // Zip Function
-    transform(make_zip_iterator(h_data0.begin(), h_data1.begin(), h_data2.begin()),
-              make_zip_iterator(h_data0.end(), h_data1.end(), h_data2.end()),
-              h_result_zip.begin(),
-              make_zip_function(SumThree{}));
-    transform(make_zip_iterator(d_data0.begin(), d_data1.begin(), d_data2.begin()),
-              make_zip_iterator(d_data0.end(), d_data1.end(), d_data2.end()),
-              d_result_zip.begin(),
-              make_zip_function(SumThree{}));
+    thrust::transform(make_zip_iterator(h_data0.begin(), h_data1.begin(), h_data2.begin()),
+                      make_zip_iterator(h_data0.end(), h_data1.end(), h_data2.end()),
+                      h_result_zip.begin(),
+                      make_zip_function(SumThree{}));
+    thrust::transform(make_zip_iterator(d_data0.begin(), d_data1.begin(), d_data2.begin()),
+                      make_zip_iterator(d_data0.end(), d_data1.end(), d_data2.end()),
+                      d_result_zip.begin(),
+                      make_zip_function(SumThree{}));
 
     ASSERT_EQUAL(h_result_tuple, h_result_zip);
     ASSERT_EQUAL(h_result_tuple, d_result_zip);

From b08d6fa54418681e995ce9e51ced8b58f3a0f723 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Tue, 25 Feb 2025 19:45:35 +0100
Subject: [PATCH 09/17] Drop/internalize some macros (#3936)

Drop: THRUST_HOST_BACKEND
Internalize: THRUST_TUNING_ARCH and CUDA_CUB_RET_IF_FAIL
---
 thrust/thrust/detail/config/host_system.h     |  4 --
 thrust/thrust/system/cuda/detail/copy_if.h    |  8 ++--
 thrust/thrust/system/cuda/detail/core/util.h  | 42 +++++++++----------
 thrust/thrust/system/cuda/detail/extrema.h    | 16 +++----
 thrust/thrust/system/cuda/detail/partition.h  |  8 ++--
 thrust/thrust/system/cuda/detail/reduce.h     | 16 +++----
 .../thrust/system/cuda/detail/reduce_by_key.h | 10 ++---
 .../system/cuda/detail/set_operations.h       | 12 +++---
 thrust/thrust/system/cuda/detail/unique.h     | 10 ++---
 .../thrust/system/cuda/detail/unique_by_key.h |  8 ++--
 10 files changed, 63 insertions(+), 71 deletions(-)

diff --git a/thrust/thrust/detail/config/host_system.h b/thrust/thrust/detail/config/host_system.h
index 8fc6a4f9c1c..5c61440446d 100644
--- a/thrust/thrust/detail/config/host_system.h
+++ b/thrust/thrust/detail/config/host_system.h
@@ -35,10 +35,6 @@
 #  define THRUST_HOST_SYSTEM THRUST_HOST_SYSTEM_CPP
 #endif // THRUST_HOST_SYSTEM
 
-#ifdef THRUST_HOST_BACKEND
-#  error THRUST_HOST_BACKEND is no longer supported; use THRUST_HOST_SYSTEM instead.
-#endif // THRUST_HOST_BACKEND
-
 #if THRUST_HOST_SYSTEM == THRUST_HOST_SYSTEM_CPP
 #  define __THRUST_HOST_SYSTEM_NAMESPACE cpp
 #elif THRUST_HOST_SYSTEM == THRUST_HOST_SYSTEM_OMP
diff --git a/thrust/thrust/system/cuda/detail/copy_if.h b/thrust/thrust/system/cuda/detail/copy_if.h
index 33ba756eafa..94dd441c15b 100644
--- a/thrust/thrust/system/cuda/detail/copy_if.h
+++ b/thrust/thrust/system/cuda/detail/copy_if.h
@@ -129,10 +129,10 @@ struct DispatchCopyIf
           equality_op_t{},
           num_items,
           stream);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     status = cub::detail::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     // Return if we're only querying temporary storage requirements
     if (d_temp_storage == nullptr)
@@ -163,11 +163,11 @@ struct DispatchCopyIf
           equality_op_t{},
           num_items,
           stream);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     // Get number of selected items
     status = cuda_cub::synchronize(policy);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
     OffsetT num_selected = get_value(policy, d_num_selected_out);
     thrust::advance(output, num_selected);
     return status;
diff --git a/thrust/thrust/system/cuda/detail/core/util.h b/thrust/thrust/system/cuda/detail/core/util.h
index a5f87026d7d..12eff7a94a8 100644
--- a/thrust/thrust/system/cuda/detail/core/util.h
+++ b/thrust/thrust/system/cuda/detail/core/util.h
@@ -59,25 +59,6 @@ namespace cuda_cub
 {
 namespace core
 {
-
-#ifdef _NVHPC_CUDA
-#  if (__NVCOMPILER_CUDA_ARCH__ >= 600)
-// deprecated [since 2.8]
-#    define THRUST_TUNING_ARCH sm60
-#  else
-// deprecated [since 2.8]
-#    define THRUST_TUNING_ARCH sm52
-#  endif
-#else
-#  if (__CUDA_ARCH__ >= 600)
-// deprecated [since 2.8]
-#    define THRUST_TUNING_ARCH sm60
-#  else
-// deprecated [since 2.8]
-#    define THRUST_TUNING_ARCH sm52
-#  endif
-#endif
-
 namespace detail
 {
 /// Typelist - a container of types
@@ -176,7 +157,21 @@ struct specialize_plan_impl_match<P, typelist<SM, SMs...>>
     : ::cuda::std::conditional<has_sm_tuning<P, SM>::value, P<SM>, specialize_plan_impl_match<P, typelist<SMs...>>>::type
 {};
 
-template <template <class> class Plan, class SM = THRUST_TUNING_ARCH>
+#ifdef _NVHPC_CUDA
+#  if (__NVCOMPILER_CUDA_ARCH__ >= 600)
+#    define _THRUST_TUNING_ARCH sm60
+#  else
+#    define _THRUST_TUNING_ARCH sm52
+#  endif
+#else
+#  if (__CUDA_ARCH__ >= 600)
+#    define _THRUST_TUNING_ARCH sm60
+#  else
+#    define _THRUST_TUNING_ARCH sm52
+#  endif
+#endif
+
+template <template <class> class Plan, class SM = _THRUST_TUNING_ARCH>
 struct specialize_plan_msvc10_war
 {
   // if Plan has tuning type, this means it has SM-specific tuning
@@ -187,10 +182,12 @@ struct specialize_plan_msvc10_war
                                         Plan<SM>>;
 };
 
-template <template <class> class Plan, class SM = THRUST_TUNING_ARCH>
+template <template <class> class Plan, class SM = _THRUST_TUNING_ARCH>
 struct specialize_plan : specialize_plan_msvc10_war<Plan, SM>::type::type
 {};
 
+#undef _THRUST_TUNING_ARCH
+
 /////////////////////////
 /////////////////////////
 /////////////////////////
@@ -585,8 +582,7 @@ THRUST_RUNTIME_FUNCTION inline int get_ptx_version()
   return ptx_version;
 }
 
-// Deprecated [Since 2.8]
-#define CUDA_CUB_RET_IF_FAIL(e)                \
+#define _CUDA_CUB_RET_IF_FAIL(e)               \
   {                                            \
     auto const error = (e);                    \
     if (cub::Debug(error, __FILE__, __LINE__)) \
diff --git a/thrust/thrust/system/cuda/detail/extrema.h b/thrust/thrust/system/cuda/detail/extrema.h
index 74891347cd1..60a093c05dc 100644
--- a/thrust/thrust/system/cuda/detail/extrema.h
+++ b/thrust/thrust/system/cuda/detail/extrema.h
@@ -216,13 +216,13 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
 
     reduce_agent ra(reduce_plan, num_items, stream, vshmem_ptr, "reduce_agent: single_tile only");
     ra.launch(input_it, output_it, num_items, reduction_op);
-    CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+    _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
   }
   else
   {
     // regular size
     cuda_optional<int> sm_count = core::detail::get_sm_count();
-    CUDA_CUB_RET_IF_FAIL(sm_count.status());
+    _CUDA_CUB_RET_IF_FAIL(sm_count.status());
 
     // reduction will not use more cta counts than requested
     cuda_optional<int> max_blocks_per_sm = reduce_agent::template get_max_blocks_per_sm<
@@ -232,7 +232,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
       cub::GridEvenShare<Size>,
       cub::GridQueue<UnsignedSize>,
       ReductionOp>(reduce_plan);
-    CUDA_CUB_RET_IF_FAIL(max_blocks_per_sm.status());
+    _CUDA_CUB_RET_IF_FAIL(max_blocks_per_sm.status());
 
     int reduce_device_occupancy = (int) max_blocks_per_sm * sm_count;
 
@@ -255,7 +255,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
       vshmem_size // size of virtualized shared memory storage
     };
     status = cub::detail::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
     if (d_temp_storage == nullptr)
     {
       return status;
@@ -286,17 +286,17 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
       drain_plan.grid_size = 1;
       drain_agent da(drain_plan, stream, "__reduce::drain_agent");
       da.launch(queue, num_items);
-      CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+      _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
     }
     else
     {
-      CUDA_CUB_RET_IF_FAIL(cudaErrorNotSupported);
+      _CUDA_CUB_RET_IF_FAIL(cudaErrorNotSupported);
     }
 
     reduce_plan.grid_size = reduce_grid_size;
     reduce_agent ra(reduce_plan, stream, vshmem_ptr, "reduce_agent: regular size reduce");
     ra.launch(input_it, d_block_reductions, num_items, even_share, queue, reduction_op);
-    CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+    _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
 
     using reduce_agent_single = AgentLauncher<__reduce::ReduceAgent<T*, OutputIt, T, Size, ReductionOp>>;
 
@@ -304,7 +304,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
     reduce_agent_single ra1(reduce_plan, stream, vshmem_ptr, "reduce_agent: single tile reduce");
 
     ra1.launch(d_block_reductions, output_it, reduce_grid_size, reduction_op);
-    CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+    _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
   }
 
   return status;
diff --git a/thrust/thrust/system/cuda/detail/partition.h b/thrust/thrust/system/cuda/detail/partition.h
index 690e1e4dd6d..f0ce7a1c62d 100644
--- a/thrust/thrust/system/cuda/detail/partition.h
+++ b/thrust/thrust/system/cuda/detail/partition.h
@@ -106,10 +106,10 @@ struct DispatchPartitionIf
                                             equality_op_t{},
                                             num_items,
                                             stream);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     status = cub::detail::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     // Return if we're only querying temporary storage requirements
     if (d_temp_storage == nullptr)
@@ -146,11 +146,11 @@ struct DispatchPartitionIf
                                             equality_op_t{},
                                             num_items,
                                             stream);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     // Get number of selected items
     status = cuda_cub::synchronize(policy);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
     num_selected = static_cast<std::size_t>(get_value(policy, d_num_selected_out));
 
     return status;
diff --git a/thrust/thrust/system/cuda/detail/reduce.h b/thrust/thrust/system/cuda/detail/reduce.h
index 0c266f71b54..dabafcaaf49 100644
--- a/thrust/thrust/system/cuda/detail/reduce.h
+++ b/thrust/thrust/system/cuda/detail/reduce.h
@@ -639,13 +639,13 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
 
     reduce_agent ra(reduce_plan, num_items, stream, vshmem_ptr, "reduce_agent: single_tile only");
     ra.launch(input_it, output_it, num_items, reduction_op, init);
-    CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+    _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
   }
   else
   {
     // regular size
     cuda_optional<int> sm_count = core::detail::get_sm_count();
-    CUDA_CUB_RET_IF_FAIL(sm_count.status());
+    _CUDA_CUB_RET_IF_FAIL(sm_count.status());
 
     // reduction will not use more cta counts than requested
     cuda_optional<int> max_blocks_per_sm = reduce_agent::template get_max_blocks_per_sm<
@@ -655,7 +655,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
       cub::GridEvenShare<Size>,
       cub::GridQueue<UnsignedSize>,
       ReductionOp>(reduce_plan);
-    CUDA_CUB_RET_IF_FAIL(max_blocks_per_sm.status());
+    _CUDA_CUB_RET_IF_FAIL(max_blocks_per_sm.status());
 
     int reduce_device_occupancy = (int) max_blocks_per_sm * sm_count;
 
@@ -678,7 +678,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
       vshmem_size // size of virtualized shared memory storage
     };
     status = cub::detail::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
     if (d_temp_storage == nullptr)
     {
       return status;
@@ -709,17 +709,17 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
       drain_plan.grid_size = 1;
       drain_agent da(drain_plan, stream, "__reduce::drain_agent");
       da.launch(queue, num_items);
-      CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+      _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
     }
     else
     {
-      CUDA_CUB_RET_IF_FAIL(cudaErrorNotSupported);
+      _CUDA_CUB_RET_IF_FAIL(cudaErrorNotSupported);
     }
 
     reduce_plan.grid_size = reduce_grid_size;
     reduce_agent ra(reduce_plan, stream, vshmem_ptr, "reduce_agent: regular size reduce");
     ra.launch(input_it, d_block_reductions, num_items, even_share, queue, reduction_op);
-    CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+    _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
 
     using reduce_agent_single = AgentLauncher<ReduceAgent<T*, OutputIt, T, Size, ReductionOp>>;
 
@@ -727,7 +727,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
     reduce_agent_single ra1(reduce_plan, stream, vshmem_ptr, "reduce_agent: single tile reduce");
 
     ra1.launch(d_block_reductions, output_it, reduce_grid_size, reduction_op, init);
-    CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+    _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
   }
 
   return status;
diff --git a/thrust/thrust/system/cuda/detail/reduce_by_key.h b/thrust/thrust/system/cuda/detail/reduce_by_key.h
index 902d5574d91..1982ab6c19e 100644
--- a/thrust/thrust/system/cuda/detail/reduce_by_key.h
+++ b/thrust/thrust/system/cuda/detail/reduce_by_key.h
@@ -758,11 +758,11 @@ THRUST_RUNTIME_FUNCTION cudaError_t doit_step(
 
   size_t allocation_sizes[2] = {9, vshmem_size};
   status                     = ScanTileState::AllocationSize(static_cast<int>(num_tiles), allocation_sizes[0]);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   void* allocations[2] = {nullptr, nullptr};
   status = cub::detail::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   if (d_temp_storage == nullptr)
   {
@@ -771,11 +771,11 @@ THRUST_RUNTIME_FUNCTION cudaError_t doit_step(
 
   ScanTileState tile_state;
   status = tile_state.Init(static_cast<int>(num_tiles), allocations[0], allocation_sizes[0]);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   init_agent ia(init_plan, num_tiles, stream, "reduce_by_key::init_agent");
   ia.launch(tile_state, num_tiles, num_runs_output_it);
-  CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+  _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
 
   char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[1] : nullptr;
 
@@ -791,7 +791,7 @@ THRUST_RUNTIME_FUNCTION cudaError_t doit_step(
     reduction_op,
     num_items,
     num_tiles);
-  CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+  _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
   return status;
 }
 
diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h
index 85d03ae51cb..b1344061495 100644
--- a/thrust/thrust/system/cuda/detail/set_operations.h
+++ b/thrust/thrust/system/cuda/detail/set_operations.h
@@ -1075,7 +1075,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
 
   size_t tile_agent_storage;
   status = ScanTileState::AllocationSize(static_cast<int>(num_tiles), tile_agent_storage);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   size_t vshmem_storage          = core::detail::vshmem_size(set_op_plan.shared_memory_size, num_tiles);
   size_t partition_agent_storage = (num_tiles + 1) * sizeof(Size) * 2;
@@ -1084,7 +1084,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
   size_t allocation_sizes[3] = {tile_agent_storage, partition_agent_storage, vshmem_storage};
 
   status = core::detail::alias_storage(d_temp_storage, temp_storage_size, allocations, allocation_sizes);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   if (d_temp_storage == nullptr)
   {
@@ -1093,18 +1093,18 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
 
   ScanTileState tile_state;
   status = tile_state.Init(static_cast<int>(num_tiles), allocations[0], allocation_sizes[0]);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   pair<Size, Size>* partitions = (pair<Size, Size>*) allocations[1];
   char* vshmem_ptr             = vshmem_storage > 0 ? (char*) allocations[2] : nullptr;
 
   init_agent ia(init_plan, num_tiles, stream, "set_op::init_agent");
   ia.launch(tile_state, num_tiles);
-  CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+  _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
 
   partition_agent pa(partition_plan, num_tiles + 1, stream, "set_op::partition agent");
   pa.launch(keys1, keys2, num_keys1, num_keys2, num_tiles + 1, partitions, compare_op, tile_size);
-  CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+  _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
 
   set_op_agent sa(set_op_plan, keys_total, stream, vshmem_ptr, "set_op::set_op_agent");
   sa.launch(
@@ -1121,7 +1121,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
     partitions,
     output_count,
     tile_state);
-  CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+  _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
 
   return status;
 }
diff --git a/thrust/thrust/system/cuda/detail/unique.h b/thrust/thrust/system/cuda/detail/unique.h
index 7818aa7f916..216e7996565 100644
--- a/thrust/thrust/system/cuda/detail/unique.h
+++ b/thrust/thrust/system/cuda/detail/unique.h
@@ -482,12 +482,12 @@ static cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
   cudaError_t status         = cudaSuccess;
   size_t allocation_sizes[2] = {0, vshmem_size};
   status                     = ScanTileState::AllocationSize(static_cast<int>(num_tiles), allocation_sizes[0]);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   void* allocations[2] = {nullptr, nullptr};
   //
   status = cub::detail::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   if (d_temp_storage == nullptr)
   {
@@ -496,12 +496,12 @@ static cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
 
   ScanTileState tile_status;
   status = tile_status.Init(static_cast<int>(num_tiles), allocations[0], allocation_sizes[0]);
-  CUDA_CUB_RET_IF_FAIL(status);
+  _CUDA_CUB_RET_IF_FAIL(status);
 
   num_tiles = ::cuda::std::max<size_t>(1, num_tiles);
   init_agent ia(init_plan, num_tiles, stream, "unique_by_key::init_agent");
   ia.launch(tile_status, num_tiles, num_selected_out);
-  CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+  _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
 
   if (num_items == 0)
   {
@@ -512,7 +512,7 @@ static cudaError_t THRUST_RUNTIME_FUNCTION doit_step(
 
   unique_agent ua(unique_plan, num_items, stream, vshmem_ptr, "unique_by_key::unique_agent");
   ua.launch(items_in, items_out, binary_pred, num_selected_out, num_items, tile_status, num_tiles);
-  CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
+  _CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError());
   return status;
 }
 
diff --git a/thrust/thrust/system/cuda/detail/unique_by_key.h b/thrust/thrust/system/cuda/detail/unique_by_key.h
index 468c1bdae6b..6edb23cee86 100644
--- a/thrust/thrust/system/cuda/detail/unique_by_key.h
+++ b/thrust/thrust/system/cuda/detail/unique_by_key.h
@@ -122,10 +122,10 @@ struct DispatchUniqueByKey
       num_items,
       binary_pred,
       stream);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     status = cub::detail::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     // Return if we're only querying temporary storage requirements
     if (d_temp_storage == nullptr)
@@ -155,11 +155,11 @@ struct DispatchUniqueByKey
       num_items,
       binary_pred,
       stream);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
 
     // Get number of selected items
     status = cuda_cub::synchronize(policy);
-    CUDA_CUB_RET_IF_FAIL(status);
+    _CUDA_CUB_RET_IF_FAIL(status);
     OffsetT num_selected = get_value(policy, d_num_selected_out);
 
     result_end = thrust::make_pair(keys_out + num_selected, values_out + num_selected);

From 512f674fc315bce72baa1da45972945d4dd2120a Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Tue, 25 Feb 2025 19:48:14 +0100
Subject: [PATCH 10/17] Drop public access to RegBoundScaling/MemBoundScaling
 (#3934)

---
 cub/cub/util_arch.cuh | 9 ---------
 1 file changed, 9 deletions(-)

diff --git a/cub/cub/util_arch.cuh b/cub/cub/util_arch.cuh
index 16552e12258..82228215b7d 100644
--- a/cub/cub/util_arch.cuh
+++ b/cub/cub/util_arch.cuh
@@ -137,15 +137,6 @@ struct MemBoundScaling
 };
 
 } // namespace detail
-
-template <int Nominal4ByteBlockThreads, int Nominal4ByteItemsPerThread, typename T>
-using RegBoundScaling CCCL_DEPRECATED_BECAUSE("Internal implementation detail") =
-  detail::RegBoundScaling<Nominal4ByteBlockThreads, Nominal4ByteItemsPerThread, T>;
-
-template <int Nominal4ByteBlockThreads, int Nominal4ByteItemsPerThread, typename T>
-using MemBoundScaling CCCL_DEPRECATED_BECAUSE("Internal implementation detail") =
-  detail::RegBoundScaling<Nominal4ByteBlockThreads, Nominal4ByteItemsPerThread, T>;
-
 #endif // Do not document
 
 CUB_NAMESPACE_END

From 1e2f3248181a87b716c41c9d61c38489d62303d0 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Tue, 25 Feb 2025 23:22:48 +0100
Subject: [PATCH 11/17] [PTX] Add shl,shr, bmsk, prmt (#3939)

---
 docs/libcudacxx/ptx/instructions.rst          |  12 +-
 docs/libcudacxx/ptx/instructions/bmsk.rst     |   9 +
 .../ptx/instructions/generated/bmsk.rst       |  22 ++
 .../ptx/instructions/generated/prmt.rst       |  79 ++++++
 .../ptx/instructions/generated/shl.rst        |  32 +++
 .../ptx/instructions/generated/shr.rst        |  62 +++++
 docs/libcudacxx/ptx/instructions/prmt.rst     |   9 +
 docs/libcudacxx/ptx/instructions/shl.rst      |   9 +
 docs/libcudacxx/ptx/instructions/shr.rst      |   9 +
 libcudacxx/codegen/add_ptx_instruction.py     |  28 ++-
 .../include/cuda/__ptx/instructions/bmsk.h    |  37 +++
 .../cuda/__ptx/instructions/generated/bmsk.h  |  54 ++++
 .../cuda/__ptx/instructions/generated/prmt.h  | 230 ++++++++++++++++++
 .../cuda/__ptx/instructions/generated/shl.h   |  96 ++++++++
 .../cuda/__ptx/instructions/generated/shr.h   | 168 +++++++++++++
 .../include/cuda/__ptx/instructions/prmt.h    |  37 +++
 .../include/cuda/__ptx/instructions/shl.h     |  37 +++
 .../include/cuda/__ptx/instructions/shr.h     |  37 +++
 libcudacxx/include/cuda/ptx                   |   4 +
 .../test/libcudacxx/cuda/ptx/generated/bmsk.h |  34 +++
 .../test/libcudacxx/cuda/ptx/generated/prmt.h |  74 ++++++
 .../test/libcudacxx/cuda/ptx/generated/shl.h  |  39 +++
 .../test/libcudacxx/cuda/ptx/generated/shr.h  |  60 +++++
 .../cuda/ptx/ptx.bmsk.compile.pass.cpp        |  22 ++
 .../cuda/ptx/ptx.prmt.compile.pass.cpp        |  22 ++
 .../cuda/ptx/ptx.shl.compile.pass.cpp         |  22 ++
 .../cuda/ptx/ptx.shr.compile.pass.cpp         |  22 ++
 27 files changed, 1257 insertions(+), 9 deletions(-)
 create mode 100644 docs/libcudacxx/ptx/instructions/bmsk.rst
 create mode 100644 docs/libcudacxx/ptx/instructions/generated/bmsk.rst
 create mode 100644 docs/libcudacxx/ptx/instructions/generated/prmt.rst
 create mode 100644 docs/libcudacxx/ptx/instructions/generated/shl.rst
 create mode 100644 docs/libcudacxx/ptx/instructions/generated/shr.rst
 create mode 100644 docs/libcudacxx/ptx/instructions/prmt.rst
 create mode 100644 docs/libcudacxx/ptx/instructions/shl.rst
 create mode 100644 docs/libcudacxx/ptx/instructions/shr.rst
 create mode 100644 libcudacxx/include/cuda/__ptx/instructions/bmsk.h
 create mode 100644 libcudacxx/include/cuda/__ptx/instructions/generated/bmsk.h
 create mode 100644 libcudacxx/include/cuda/__ptx/instructions/generated/prmt.h
 create mode 100644 libcudacxx/include/cuda/__ptx/instructions/generated/shl.h
 create mode 100644 libcudacxx/include/cuda/__ptx/instructions/generated/shr.h
 create mode 100644 libcudacxx/include/cuda/__ptx/instructions/prmt.h
 create mode 100644 libcudacxx/include/cuda/__ptx/instructions/shl.h
 create mode 100644 libcudacxx/include/cuda/__ptx/instructions/shr.h
 create mode 100644 libcudacxx/test/libcudacxx/cuda/ptx/generated/bmsk.h
 create mode 100644 libcudacxx/test/libcudacxx/cuda/ptx/generated/prmt.h
 create mode 100644 libcudacxx/test/libcudacxx/cuda/ptx/generated/shl.h
 create mode 100644 libcudacxx/test/libcudacxx/cuda/ptx/generated/shr.h
 create mode 100644 libcudacxx/test/libcudacxx/cuda/ptx/ptx.bmsk.compile.pass.cpp
 create mode 100644 libcudacxx/test/libcudacxx/cuda/ptx/ptx.prmt.compile.pass.cpp
 create mode 100644 libcudacxx/test/libcudacxx/cuda/ptx/ptx.shl.compile.pass.cpp
 create mode 100644 libcudacxx/test/libcudacxx/cuda/ptx/ptx.shr.compile.pass.cpp

diff --git a/docs/libcudacxx/ptx/instructions.rst b/docs/libcudacxx/ptx/instructions.rst
index 6225bf2c247..373e1afbc89 100644
--- a/docs/libcudacxx/ptx/instructions.rst
+++ b/docs/libcudacxx/ptx/instructions.rst
@@ -6,6 +6,10 @@ PTX Instructions
 .. toctree::
    :maxdepth: 1
 
+   instructions/shr
+   instructions/shl
+   instructions/bmsk
+   instructions/prmt
    instructions/barrier_cluster
    instructions/bfind
    instructions/clusterlaunchcontrol
@@ -88,7 +92,7 @@ Instructions by section
    * - `szext <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-szext>`__
      - No
    * - `bmsk <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-bmsk>`__
-     - No
+     - Yes, CCCL 3.0.0 / CUDA 13.0
    * - `dp4a <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-dp4a>`__
      - No
    * - `dp2a <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-dp2a>`__
@@ -238,9 +242,9 @@ Instructions by section
    * - `shf <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-shf>`__
      - No
    * - `shl <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-shl>`__
-     - No
+     - Yes, CCCL 3.0.0 / CUDA 13.0
    * - `shr <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-shr>`__
-     - No
+     - Yes, CCCL 3.0.0 / CUDA 13.0
 
 .. list-table:: `Data Movement and Conversion Instructions <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions>`__
    :widths: 50 50
@@ -255,7 +259,7 @@ Instructions by section
    * - `shfl.sync <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-shfl-sync>`__
      - Yes, CCCL 2.9.0 / CUDA 12.9
    * - `prmt <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-prmt>`__
-     - No
+     - Yes, CCCL 3.0.0 / CUDA 13.0
    * - `ld <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld>`__
      - No
    * - `ld.global.nc <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-ld-global-nc>`__
diff --git a/docs/libcudacxx/ptx/instructions/bmsk.rst b/docs/libcudacxx/ptx/instructions/bmsk.rst
new file mode 100644
index 00000000000..167a3b2448d
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/bmsk.rst
@@ -0,0 +1,9 @@
+.. _libcudacxx-ptx-instructions-bmsk:
+
+bmsk
+====
+
+-  PTX ISA:
+   `bmsk <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#integer-arithmetic-instructions-bmsk>`__
+
+.. include:: generated/bmsk.rst
diff --git a/docs/libcudacxx/ptx/instructions/generated/bmsk.rst b/docs/libcudacxx/ptx/instructions/generated/bmsk.rst
new file mode 100644
index 00000000000..e75e0ccdaeb
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/generated/bmsk.rst
@@ -0,0 +1,22 @@
+..
+   This file was automatically generated. Do not edit.
+
+bmsk.clamp.b32
+^^^^^^^^^^^^^^
+.. code:: cuda
+
+   // bmsk.clamp.b32 dest, a_reg, b_reg; // PTX ISA 76, SM_70
+   template <typename = void>
+   __device__ static inline uint32_t bmsk_clamp(
+     uint32_t a_reg,
+     uint32_t b_reg);
+
+bmsk.wrap.b32
+^^^^^^^^^^^^^
+.. code:: cuda
+
+   // bmsk.wrap.b32 dest, a_reg, b_reg; // PTX ISA 76, SM_70
+   template <typename = void>
+   __device__ static inline uint32_t bmsk_wrap(
+     uint32_t a_reg,
+     uint32_t b_reg);
diff --git a/docs/libcudacxx/ptx/instructions/generated/prmt.rst b/docs/libcudacxx/ptx/instructions/generated/prmt.rst
new file mode 100644
index 00000000000..36a7578907e
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/generated/prmt.rst
@@ -0,0 +1,79 @@
+..
+   This file was automatically generated. Do not edit.
+
+prmt.b32
+^^^^^^^^
+.. code:: cuda
+
+   // prmt.b32 dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline uint32_t prmt(
+     B32 a_reg,
+     B32 b_reg,
+     uint32_t c_reg);
+
+prmt.b32.f4e
+^^^^^^^^^^^^
+.. code:: cuda
+
+   // prmt.b32.f4e dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline uint32_t prmt_f4e(
+     B32 a_reg,
+     B32 b_reg,
+     uint32_t c_reg);
+
+prmt.b32.b4e
+^^^^^^^^^^^^
+.. code:: cuda
+
+   // prmt.b32.b4e dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline uint32_t prmt_b4e(
+     B32 a_reg,
+     B32 b_reg,
+     uint32_t c_reg);
+
+prmt.b32.rc8
+^^^^^^^^^^^^
+.. code:: cuda
+
+   // prmt.b32.rc8 dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline uint32_t prmt_rc8(
+     B32 a_reg,
+     B32 b_reg,
+     uint32_t c_reg);
+
+prmt.b32.ecl
+^^^^^^^^^^^^
+.. code:: cuda
+
+   // prmt.b32.ecl dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline uint32_t prmt_ecl(
+     B32 a_reg,
+     B32 b_reg,
+     uint32_t c_reg);
+
+prmt.b32.ecr
+^^^^^^^^^^^^
+.. code:: cuda
+
+   // prmt.b32.ecr dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline uint32_t prmt_ecr(
+     B32 a_reg,
+     B32 b_reg,
+     uint32_t c_reg);
+
+prmt.b32.rc16
+^^^^^^^^^^^^^
+.. code:: cuda
+
+   // prmt.b32.rc16 dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline uint32_t prmt_rc16(
+     B32 a_reg,
+     B32 b_reg,
+     uint32_t c_reg);
diff --git a/docs/libcudacxx/ptx/instructions/generated/shl.rst b/docs/libcudacxx/ptx/instructions/generated/shl.rst
new file mode 100644
index 00000000000..8c1bca789fe
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/generated/shl.rst
@@ -0,0 +1,32 @@
+..
+   This file was automatically generated. Do not edit.
+
+shl.b16
+^^^^^^^
+.. code:: cuda
+
+   // shl.b16 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename B16, enable_if_t<sizeof(B16) == 2, bool> = true>
+   __device__ static inline B16 shl(
+     B16 a_reg,
+     uint32_t b_reg);
+
+shl.b32
+^^^^^^^
+.. code:: cuda
+
+   // shl.b32 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline B32 shl(
+     B32 a_reg,
+     uint32_t b_reg);
+
+shl.b64
+^^^^^^^
+.. code:: cuda
+
+   // shl.b64 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename B64, enable_if_t<sizeof(B64) == 8, bool> = true>
+   __device__ static inline B64 shl(
+     B64 a_reg,
+     uint32_t b_reg);
diff --git a/docs/libcudacxx/ptx/instructions/generated/shr.rst b/docs/libcudacxx/ptx/instructions/generated/shr.rst
new file mode 100644
index 00000000000..a1c5c2c74b4
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/generated/shr.rst
@@ -0,0 +1,62 @@
+..
+   This file was automatically generated. Do not edit.
+
+shr.b16
+^^^^^^^
+.. code:: cuda
+
+   // shr.b16 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename B16, enable_if_t<sizeof(B16) == 2, bool> = true>
+   __device__ static inline B16 shr(
+     B16 a_reg,
+     uint32_t b_reg);
+
+shr.b32
+^^^^^^^
+.. code:: cuda
+
+   // shr.b32 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+   __device__ static inline B32 shr(
+     B32 a_reg,
+     uint32_t b_reg);
+
+shr.b64
+^^^^^^^
+.. code:: cuda
+
+   // shr.b64 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename B64, enable_if_t<sizeof(B64) == 8, bool> = true>
+   __device__ static inline B64 shr(
+     B64 a_reg,
+     uint32_t b_reg);
+
+shr.s16
+^^^^^^^
+.. code:: cuda
+
+   // shr.s16 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename = void>
+   __device__ static inline int16_t shr(
+     int16_t a_reg,
+     uint32_t b_reg);
+
+shr.s32
+^^^^^^^
+.. code:: cuda
+
+   // shr.s32 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename = void>
+   __device__ static inline int32_t shr(
+     int32_t a_reg,
+     uint32_t b_reg);
+
+shr.s64
+^^^^^^^
+.. code:: cuda
+
+   // shr.s64 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+   template <typename = void>
+   __device__ static inline int64_t shr(
+     int64_t a_reg,
+     uint32_t b_reg);
diff --git a/docs/libcudacxx/ptx/instructions/prmt.rst b/docs/libcudacxx/ptx/instructions/prmt.rst
new file mode 100644
index 00000000000..3006a56731e
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/prmt.rst
@@ -0,0 +1,9 @@
+.. _libcudacxx-ptx-instructions-prmt:
+
+prmt
+====
+
+-  PTX ISA:
+   `prmt <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-prmt>`__
+
+.. include:: generated/prmt.rst
diff --git a/docs/libcudacxx/ptx/instructions/shl.rst b/docs/libcudacxx/ptx/instructions/shl.rst
new file mode 100644
index 00000000000..da23c233fb4
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/shl.rst
@@ -0,0 +1,9 @@
+.. _libcudacxx-ptx-instructions-shl:
+
+shl
+===
+
+-  PTX ISA:
+   `shl <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-shl>`__
+
+.. include:: generated/shl.rst
diff --git a/docs/libcudacxx/ptx/instructions/shr.rst b/docs/libcudacxx/ptx/instructions/shr.rst
new file mode 100644
index 00000000000..11922e48954
--- /dev/null
+++ b/docs/libcudacxx/ptx/instructions/shr.rst
@@ -0,0 +1,9 @@
+.. _libcudacxx-ptx-instructions-shr:
+
+shr
+===
+
+-  PTX ISA:
+   `shr <https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#logic-and-shift-instructions-shr>`__
+
+.. include:: generated/shr.rst
diff --git a/libcudacxx/codegen/add_ptx_instruction.py b/libcudacxx/codegen/add_ptx_instruction.py
index 811c9cdb021..7541d05653c 100755
--- a/libcudacxx/codegen/add_ptx_instruction.py
+++ b/libcudacxx/codegen/add_ptx_instruction.py
@@ -6,6 +6,7 @@
 test = Path("libcudacxx/test/libcudacxx/cuda/ptx")
 src = Path("libcudacxx/include/cuda/__ptx/instructions")
 ptx_header = Path("libcudacxx/include/cuda/ptx")
+instr_docs = Path("docs/libcudacxx/ptx/instructions.rst")
 
 
 def add_docs(ptx_instr, url):
@@ -28,7 +29,8 @@ def add_docs(ptx_instr, url):
 
 def add_test(ptx_instr):
     cpp_instr = ptx_instr.replace(".", "_")
-    (test / f"ptx.{ptx_instr}.compile.pass.cpp").write_text(
+    dst = test / f"ptx.{ptx_instr}.compile.pass.cpp"
+    dst.write_text(
         f"""//===----------------------------------------------------------------------===//
 //
 // Part of libcu++, the C++ Standard Library for your entire system,
@@ -99,11 +101,26 @@ def add_src(ptx_instr):
     )
 
 
-def add_include_reminder(ptx_instr):
+def add_ptx_header_include(ptx_instr):
     cpp_instr = ptx_instr.replace(".", "_")
     txt = ptx_header.read_text()
-    reminder = f"""// TODO: #include <cuda/__ptx/instructions/{cpp_instr}.h>"""
-    ptx_header.write_text(f"{txt}\n{reminder}\n")
+    # just add as first new include. clang-format will sort it in
+    idx = txt.index("#include <cuda/__ptx/instructions")
+    txt = (
+        txt[:idx]
+        + f"""#include <cuda/__ptx/instructions/{cpp_instr}.h>\n"""
+        + txt[idx:]
+    )
+    ptx_header.write_text(txt)
+
+
+def add_docs_include(ptx_instr):
+    cpp_instr = ptx_instr.replace(".", "_")
+    txt = instr_docs.read_text()
+    # just add as first new include
+    idx = txt.index("   instructions/")
+    txt = txt[:idx] + f"   instructions/{cpp_instr}\n" + txt[idx:]
+    instr_docs.write_text(txt)
 
 
 if __name__ == "__main__":
@@ -126,4 +143,5 @@ def add_include_reminder(ptx_instr):
     add_test(ptx_instr)
     add_docs(ptx_instr, url)
     add_src(ptx_instr)
-    add_include_reminder(ptx_instr)
+    add_ptx_header_include(ptx_instr)
+    add_docs_include(ptx_instr)
diff --git a/libcudacxx/include/cuda/__ptx/instructions/bmsk.h b/libcudacxx/include/cuda/__ptx/instructions/bmsk.h
new file mode 100644
index 00000000000..be5ad5a1874
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/bmsk.h
@@ -0,0 +1,37 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_PTX_BMSK_H_
+#define _CUDA_PTX_BMSK_H_
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__ptx/ptx_dot_variants.h>
+#include <cuda/__ptx/ptx_helper_functions.h>
+#include <cuda/std/cstdint>
+
+#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX
+
+#include <cuda/__ptx/instructions/generated/bmsk.h>
+
+_LIBCUDACXX_END_NAMESPACE_CUDA_PTX
+
+#endif // _CUDA_PTX_BMSK_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/bmsk.h b/libcudacxx/include/cuda/__ptx/instructions/generated/bmsk.h
new file mode 100644
index 00000000000..0a0a244e1ae
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/generated/bmsk.h
@@ -0,0 +1,54 @@
+// This file was automatically generated. Do not edit.
+
+#ifndef _CUDA_PTX_GENERATED_BMSK_H_
+#define _CUDA_PTX_GENERATED_BMSK_H_
+
+/*
+// bmsk.clamp.b32 dest, a_reg, b_reg; // PTX ISA 76, SM_70
+template <typename = void>
+__device__ static inline uint32_t bmsk_clamp(
+  uint32_t a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 760
+extern "C" _CCCL_DEVICE void __cuda_ptx_bmsk_clamp_is_not_supported_before_SM_70__();
+template <typename = void>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t bmsk_clamp(_CUDA_VSTD::uint32_t __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 700
+  _CUDA_VSTD::uint32_t __dest;
+  asm("bmsk.clamp.b32 %0, %1, %2;" : "=r"(__dest) : "r"(__a_reg), "r"(__b_reg) :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_bmsk_clamp_is_not_supported_before_SM_70__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 760
+
+/*
+// bmsk.wrap.b32 dest, a_reg, b_reg; // PTX ISA 76, SM_70
+template <typename = void>
+__device__ static inline uint32_t bmsk_wrap(
+  uint32_t a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 760
+extern "C" _CCCL_DEVICE void __cuda_ptx_bmsk_wrap_is_not_supported_before_SM_70__();
+template <typename = void>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t bmsk_wrap(_CUDA_VSTD::uint32_t __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 700
+  _CUDA_VSTD::uint32_t __dest;
+  asm("bmsk.wrap.b32 %0, %1, %2;" : "=r"(__dest) : "r"(__a_reg), "r"(__b_reg) :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_bmsk_wrap_is_not_supported_before_SM_70__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 760
+
+#endif // _CUDA_PTX_GENERATED_BMSK_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/prmt.h b/libcudacxx/include/cuda/__ptx/instructions/generated/prmt.h
new file mode 100644
index 00000000000..53a1533960d
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/generated/prmt.h
@@ -0,0 +1,230 @@
+// This file was automatically generated. Do not edit.
+
+#ifndef _CUDA_PTX_GENERATED_PRMT_H_
+#define _CUDA_PTX_GENERATED_PRMT_H_
+
+/*
+// prmt.b32 dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline uint32_t prmt(
+  B32 a_reg,
+  B32 b_reg,
+  uint32_t c_reg);
+*/
+#if __cccl_ptx_isa >= 200
+extern "C" _CCCL_DEVICE void __cuda_ptx_prmt_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t prmt(_B32 __a_reg, _B32 __b_reg, _CUDA_VSTD::uint32_t __c_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("prmt.b32 %0, %1, %2, %3;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)),
+        "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__b_reg)),
+        "r"(__c_reg)
+      :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_prmt_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 200
+
+/*
+// prmt.b32.f4e dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline uint32_t prmt_f4e(
+  B32 a_reg,
+  B32 b_reg,
+  uint32_t c_reg);
+*/
+#if __cccl_ptx_isa >= 200
+extern "C" _CCCL_DEVICE void __cuda_ptx_prmt_f4e_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t prmt_f4e(_B32 __a_reg, _B32 __b_reg, _CUDA_VSTD::uint32_t __c_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("prmt.b32.f4e %0, %1, %2, %3;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)),
+        "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__b_reg)),
+        "r"(__c_reg)
+      :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_prmt_f4e_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 200
+
+/*
+// prmt.b32.b4e dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline uint32_t prmt_b4e(
+  B32 a_reg,
+  B32 b_reg,
+  uint32_t c_reg);
+*/
+#if __cccl_ptx_isa >= 200
+extern "C" _CCCL_DEVICE void __cuda_ptx_prmt_b4e_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t prmt_b4e(_B32 __a_reg, _B32 __b_reg, _CUDA_VSTD::uint32_t __c_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("prmt.b32.b4e %0, %1, %2, %3;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)),
+        "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__b_reg)),
+        "r"(__c_reg)
+      :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_prmt_b4e_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 200
+
+/*
+// prmt.b32.rc8 dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline uint32_t prmt_rc8(
+  B32 a_reg,
+  B32 b_reg,
+  uint32_t c_reg);
+*/
+#if __cccl_ptx_isa >= 200
+extern "C" _CCCL_DEVICE void __cuda_ptx_prmt_rc8_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t prmt_rc8(_B32 __a_reg, _B32 __b_reg, _CUDA_VSTD::uint32_t __c_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("prmt.b32.rc8 %0, %1, %2, %3;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)),
+        "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__b_reg)),
+        "r"(__c_reg)
+      :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_prmt_rc8_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 200
+
+/*
+// prmt.b32.ecl dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline uint32_t prmt_ecl(
+  B32 a_reg,
+  B32 b_reg,
+  uint32_t c_reg);
+*/
+#if __cccl_ptx_isa >= 200
+extern "C" _CCCL_DEVICE void __cuda_ptx_prmt_ecl_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t prmt_ecl(_B32 __a_reg, _B32 __b_reg, _CUDA_VSTD::uint32_t __c_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("prmt.b32.ecl %0, %1, %2, %3;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)),
+        "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__b_reg)),
+        "r"(__c_reg)
+      :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_prmt_ecl_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 200
+
+/*
+// prmt.b32.ecr dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline uint32_t prmt_ecr(
+  B32 a_reg,
+  B32 b_reg,
+  uint32_t c_reg);
+*/
+#if __cccl_ptx_isa >= 200
+extern "C" _CCCL_DEVICE void __cuda_ptx_prmt_ecr_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t prmt_ecr(_B32 __a_reg, _B32 __b_reg, _CUDA_VSTD::uint32_t __c_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("prmt.b32.ecr %0, %1, %2, %3;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)),
+        "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__b_reg)),
+        "r"(__c_reg)
+      :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_prmt_ecr_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 200
+
+/*
+// prmt.b32.rc16 dest, a_reg, b_reg, c_reg; // PTX ISA 20, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline uint32_t prmt_rc16(
+  B32 a_reg,
+  B32 b_reg,
+  uint32_t c_reg);
+*/
+#if __cccl_ptx_isa >= 200
+extern "C" _CCCL_DEVICE void __cuda_ptx_prmt_rc16_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _CUDA_VSTD::uint32_t prmt_rc16(_B32 __a_reg, _B32 __b_reg, _CUDA_VSTD::uint32_t __c_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("prmt.b32.rc16 %0, %1, %2, %3;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)),
+        "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__b_reg)),
+        "r"(__c_reg)
+      :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_prmt_rc16_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 200
+
+#endif // _CUDA_PTX_GENERATED_PRMT_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/shl.h b/libcudacxx/include/cuda/__ptx/instructions/generated/shl.h
new file mode 100644
index 00000000000..9d7de603594
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/generated/shl.h
@@ -0,0 +1,96 @@
+// This file was automatically generated. Do not edit.
+
+#ifndef _CUDA_PTX_GENERATED_SHL_H_
+#define _CUDA_PTX_GENERATED_SHL_H_
+
+/*
+// shl.b16 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename B16, enable_if_t<sizeof(B16) == 2, bool> = true>
+__device__ static inline B16 shl(
+  B16 a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shl_is_not_supported_before_SM_50__();
+template <typename _B16, _CUDA_VSTD::enable_if_t<sizeof(_B16) == 2, bool> = true>
+_CCCL_DEVICE static inline _B16 shl(_B16 __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+  static_assert(sizeof(_B16) == 2, "");
+  static_assert(sizeof(_B16) == 2, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint16_t __dest;
+  asm("shl.b16 %0, %1, %2;"
+      : "=h"(__dest)
+      : "h"(/*as_b16*/ *reinterpret_cast<const _CUDA_VSTD::int16_t*>(&__a_reg)), "r"(__b_reg)
+      :);
+  return *reinterpret_cast<_B16*>(&__dest);
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shl_is_not_supported_before_SM_50__();
+  _CUDA_VSTD::uint16_t __err_out_var = 0;
+  return *reinterpret_cast<_B16*>(&__err_out_var);
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+/*
+// shl.b32 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline B32 shl(
+  B32 a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shl_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _B32 shl(_B32 __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("shl.b32 %0, %1, %2;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)), "r"(__b_reg)
+      :);
+  return *reinterpret_cast<_B32*>(&__dest);
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shl_is_not_supported_before_SM_50__();
+  _CUDA_VSTD::uint32_t __err_out_var = 0;
+  return *reinterpret_cast<_B32*>(&__err_out_var);
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+/*
+// shl.b64 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename B64, enable_if_t<sizeof(B64) == 8, bool> = true>
+__device__ static inline B64 shl(
+  B64 a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shl_is_not_supported_before_SM_50__();
+template <typename _B64, _CUDA_VSTD::enable_if_t<sizeof(_B64) == 8, bool> = true>
+_CCCL_DEVICE static inline _B64 shl(_B64 __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+  static_assert(sizeof(_B64) == 8, "");
+  static_assert(sizeof(_B64) == 8, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint64_t __dest;
+  asm("shl.b64 %0, %1, %2;"
+      : "=l"(__dest)
+      : "l"(/*as_b64*/ *reinterpret_cast<const _CUDA_VSTD::int64_t*>(&__a_reg)), "r"(__b_reg)
+      :);
+  return *reinterpret_cast<_B64*>(&__dest);
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shl_is_not_supported_before_SM_50__();
+  _CUDA_VSTD::uint64_t __err_out_var = 0;
+  return *reinterpret_cast<_B64*>(&__err_out_var);
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+#endif // _CUDA_PTX_GENERATED_SHL_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/shr.h b/libcudacxx/include/cuda/__ptx/instructions/generated/shr.h
new file mode 100644
index 00000000000..352bde59ea6
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/generated/shr.h
@@ -0,0 +1,168 @@
+// This file was automatically generated. Do not edit.
+
+#ifndef _CUDA_PTX_GENERATED_SHR_H_
+#define _CUDA_PTX_GENERATED_SHR_H_
+
+/*
+// shr.b16 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename B16, enable_if_t<sizeof(B16) == 2, bool> = true>
+__device__ static inline B16 shr(
+  B16 a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shr_is_not_supported_before_SM_50__();
+template <typename _B16, _CUDA_VSTD::enable_if_t<sizeof(_B16) == 2, bool> = true>
+_CCCL_DEVICE static inline _B16 shr(_B16 __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+  static_assert(sizeof(_B16) == 2, "");
+  static_assert(sizeof(_B16) == 2, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint16_t __dest;
+  asm("shr.b16 %0, %1, %2;"
+      : "=h"(__dest)
+      : "h"(/*as_b16*/ *reinterpret_cast<const _CUDA_VSTD::int16_t*>(&__a_reg)), "r"(__b_reg)
+      :);
+  return *reinterpret_cast<_B16*>(&__dest);
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shr_is_not_supported_before_SM_50__();
+  _CUDA_VSTD::uint16_t __err_out_var = 0;
+  return *reinterpret_cast<_B16*>(&__err_out_var);
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+/*
+// shr.b32 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename B32, enable_if_t<sizeof(B32) == 4, bool> = true>
+__device__ static inline B32 shr(
+  B32 a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shr_is_not_supported_before_SM_50__();
+template <typename _B32, _CUDA_VSTD::enable_if_t<sizeof(_B32) == 4, bool> = true>
+_CCCL_DEVICE static inline _B32 shr(_B32 __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+  static_assert(sizeof(_B32) == 4, "");
+  static_assert(sizeof(_B32) == 4, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint32_t __dest;
+  asm("shr.b32 %0, %1, %2;"
+      : "=r"(__dest)
+      : "r"(/*as_b32*/ *reinterpret_cast<const _CUDA_VSTD::int32_t*>(&__a_reg)), "r"(__b_reg)
+      :);
+  return *reinterpret_cast<_B32*>(&__dest);
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shr_is_not_supported_before_SM_50__();
+  _CUDA_VSTD::uint32_t __err_out_var = 0;
+  return *reinterpret_cast<_B32*>(&__err_out_var);
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+/*
+// shr.b64 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename B64, enable_if_t<sizeof(B64) == 8, bool> = true>
+__device__ static inline B64 shr(
+  B64 a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shr_is_not_supported_before_SM_50__();
+template <typename _B64, _CUDA_VSTD::enable_if_t<sizeof(_B64) == 8, bool> = true>
+_CCCL_DEVICE static inline _B64 shr(_B64 __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+  static_assert(sizeof(_B64) == 8, "");
+  static_assert(sizeof(_B64) == 8, "");
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::uint64_t __dest;
+  asm("shr.b64 %0, %1, %2;"
+      : "=l"(__dest)
+      : "l"(/*as_b64*/ *reinterpret_cast<const _CUDA_VSTD::int64_t*>(&__a_reg)), "r"(__b_reg)
+      :);
+  return *reinterpret_cast<_B64*>(&__dest);
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shr_is_not_supported_before_SM_50__();
+  _CUDA_VSTD::uint64_t __err_out_var = 0;
+  return *reinterpret_cast<_B64*>(&__err_out_var);
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+/*
+// shr.s16 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename = void>
+__device__ static inline int16_t shr(
+  int16_t a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shr_is_not_supported_before_SM_50__();
+template <typename = void>
+_CCCL_DEVICE static inline _CUDA_VSTD::int16_t shr(_CUDA_VSTD::int16_t __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::int16_t __dest;
+  asm("shr.s16 %0, %1, %2;" : "=h"(__dest) : "h"(__a_reg), "r"(__b_reg) :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shr_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+/*
+// shr.s32 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename = void>
+__device__ static inline int32_t shr(
+  int32_t a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shr_is_not_supported_before_SM_50__();
+template <typename = void>
+_CCCL_DEVICE static inline _CUDA_VSTD::int32_t shr(_CUDA_VSTD::int32_t __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::int32_t __dest;
+  asm("shr.s32 %0, %1, %2;" : "=r"(__dest) : "r"(__a_reg), "r"(__b_reg) :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shr_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+/*
+// shr.s64 dest, a_reg, b_reg; // PTX ISA 10, SM_50
+template <typename = void>
+__device__ static inline int64_t shr(
+  int64_t a_reg,
+  uint32_t b_reg);
+*/
+#if __cccl_ptx_isa >= 100
+extern "C" _CCCL_DEVICE void __cuda_ptx_shr_is_not_supported_before_SM_50__();
+template <typename = void>
+_CCCL_DEVICE static inline _CUDA_VSTD::int64_t shr(_CUDA_VSTD::int64_t __a_reg, _CUDA_VSTD::uint32_t __b_reg)
+{
+#  if _CCCL_CUDA_COMPILER(NVHPC) || __CUDA_ARCH__ >= 500
+  _CUDA_VSTD::int64_t __dest;
+  asm("shr.s64 %0, %1, %2;" : "=l"(__dest) : "l"(__a_reg), "r"(__b_reg) :);
+  return __dest;
+#  else
+  // Unsupported architectures will have a linker error with a semi-decent error message
+  __cuda_ptx_shr_is_not_supported_before_SM_50__();
+  return 0;
+#  endif
+}
+#endif // __cccl_ptx_isa >= 100
+
+#endif // _CUDA_PTX_GENERATED_SHR_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/prmt.h b/libcudacxx/include/cuda/__ptx/instructions/prmt.h
new file mode 100644
index 00000000000..32f9b54255c
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/prmt.h
@@ -0,0 +1,37 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_PTX_PRMT_H_
+#define _CUDA_PTX_PRMT_H_
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__ptx/ptx_dot_variants.h>
+#include <cuda/__ptx/ptx_helper_functions.h>
+#include <cuda/std/cstdint>
+
+#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX
+
+#include <cuda/__ptx/instructions/generated/prmt.h>
+
+_LIBCUDACXX_END_NAMESPACE_CUDA_PTX
+
+#endif // _CUDA_PTX_PRMT_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/shl.h b/libcudacxx/include/cuda/__ptx/instructions/shl.h
new file mode 100644
index 00000000000..4d5f96ecf04
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/shl.h
@@ -0,0 +1,37 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_PTX_SHL_H_
+#define _CUDA_PTX_SHL_H_
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__ptx/ptx_dot_variants.h>
+#include <cuda/__ptx/ptx_helper_functions.h>
+#include <cuda/std/cstdint>
+
+#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX
+
+#include <cuda/__ptx/instructions/generated/shl.h>
+
+_LIBCUDACXX_END_NAMESPACE_CUDA_PTX
+
+#endif // _CUDA_PTX_SHL_H_
diff --git a/libcudacxx/include/cuda/__ptx/instructions/shr.h b/libcudacxx/include/cuda/__ptx/instructions/shr.h
new file mode 100644
index 00000000000..7e04ddf0c3e
--- /dev/null
+++ b/libcudacxx/include/cuda/__ptx/instructions/shr.h
@@ -0,0 +1,37 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_PTX_SHR_H_
+#define _CUDA_PTX_SHR_H_
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__ptx/ptx_dot_variants.h>
+#include <cuda/__ptx/ptx_helper_functions.h>
+#include <cuda/std/cstdint>
+
+#include <nv/target> // __CUDA_MINIMUM_ARCH__ and friends
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX
+
+#include <cuda/__ptx/instructions/generated/shr.h>
+
+_LIBCUDACXX_END_NAMESPACE_CUDA_PTX
+
+#endif // _CUDA_PTX_SHR_H_
diff --git a/libcudacxx/include/cuda/ptx b/libcudacxx/include/cuda/ptx
index 7855a42ed8c..517493e66bb 100644
--- a/libcudacxx/include/cuda/ptx
+++ b/libcudacxx/include/cuda/ptx
@@ -71,6 +71,7 @@
 
 #include <cuda/__ptx/instructions/barrier_cluster.h>
 #include <cuda/__ptx/instructions/bfind.h>
+#include <cuda/__ptx/instructions/bmsk.h>
 #include <cuda/__ptx/instructions/clusterlaunchcontrol.h>
 #include <cuda/__ptx/instructions/cp_async_bulk.h>
 #include <cuda/__ptx/instructions/cp_async_bulk_commit_group.h>
@@ -90,8 +91,11 @@
 #include <cuda/__ptx/instructions/multimem_ld_reduce.h>
 #include <cuda/__ptx/instructions/multimem_red.h>
 #include <cuda/__ptx/instructions/multimem_st.h>
+#include <cuda/__ptx/instructions/prmt.h>
 #include <cuda/__ptx/instructions/red_async.h>
 #include <cuda/__ptx/instructions/shfl_sync.h>
+#include <cuda/__ptx/instructions/shl.h>
+#include <cuda/__ptx/instructions/shr.h>
 #include <cuda/__ptx/instructions/st_async.h>
 #include <cuda/__ptx/instructions/st_bulk.h>
 #include <cuda/__ptx/instructions/tcgen05_alloc.h>
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/generated/bmsk.h b/libcudacxx/test/libcudacxx/cuda/ptx/generated/bmsk.h
new file mode 100644
index 00000000000..e477bfe2951
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/generated/bmsk.h
@@ -0,0 +1,34 @@
+// This file was automatically generated. Do not edit.
+
+// We use a special strategy to force the generation of the PTX. This is mainly
+// a fight against dead-code-elimination in the NVVM layer.
+//
+// The reason we need this strategy is because certain older versions of ptxas
+// segfault when a non-sensical sequence of PTX is generated. So instead, we try
+// to force the instantiation and compilation to PTX of all the overloads of the
+// PTX wrapping functions.
+//
+// We do this by writing a function pointer of each overload to the kernel
+// parameter `fn_ptr`.
+//
+// Because `fn_ptr` is possibly visible outside this translation unit, the
+// compiler must compile all the functions which are stored.
+
+__global__ void test_bmsk(void** fn_ptr)
+{
+#if __cccl_ptx_isa >= 760
+  NV_IF_TARGET(
+    NV_PROVIDES_SM_70,
+    (
+        // bmsk.clamp.b32 dest, a_reg, b_reg;
+        * fn_ptr++ = reinterpret_cast<void*>(static_cast<uint32_t (*)(uint32_t, uint32_t)>(cuda::ptx::bmsk_clamp));));
+#endif // __cccl_ptx_isa >= 760
+
+#if __cccl_ptx_isa >= 760
+  NV_IF_TARGET(
+    NV_PROVIDES_SM_70,
+    (
+        // bmsk.wrap.b32 dest, a_reg, b_reg;
+        * fn_ptr++ = reinterpret_cast<void*>(static_cast<uint32_t (*)(uint32_t, uint32_t)>(cuda::ptx::bmsk_wrap));));
+#endif // __cccl_ptx_isa >= 760
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/generated/prmt.h b/libcudacxx/test/libcudacxx/cuda/ptx/generated/prmt.h
new file mode 100644
index 00000000000..5ae13c99a47
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/generated/prmt.h
@@ -0,0 +1,74 @@
+// This file was automatically generated. Do not edit.
+
+// We use a special strategy to force the generation of the PTX. This is mainly
+// a fight against dead-code-elimination in the NVVM layer.
+//
+// The reason we need this strategy is because certain older versions of ptxas
+// segfault when a non-sensical sequence of PTX is generated. So instead, we try
+// to force the instantiation and compilation to PTX of all the overloads of the
+// PTX wrapping functions.
+//
+// We do this by writing a function pointer of each overload to the kernel
+// parameter `fn_ptr`.
+//
+// Because `fn_ptr` is possibly visible outside this translation unit, the
+// compiler must compile all the functions which are stored.
+
+__global__ void test_prmt(void** fn_ptr)
+{
+#if __cccl_ptx_isa >= 200
+  NV_IF_TARGET(
+    NV_PROVIDES_SM_50,
+    (
+        // prmt.b32 dest, a_reg, b_reg, c_reg;
+        * fn_ptr++ = reinterpret_cast<void*>(static_cast<uint32_t (*)(int32_t, int32_t, uint32_t)>(cuda::ptx::prmt));));
+#endif // __cccl_ptx_isa >= 200
+
+#if __cccl_ptx_isa >= 200
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // prmt.b32.f4e dest, a_reg, b_reg, c_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(
+                     static_cast<uint32_t (*)(int32_t, int32_t, uint32_t)>(cuda::ptx::prmt_f4e));));
+#endif // __cccl_ptx_isa >= 200
+
+#if __cccl_ptx_isa >= 200
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // prmt.b32.b4e dest, a_reg, b_reg, c_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(
+                     static_cast<uint32_t (*)(int32_t, int32_t, uint32_t)>(cuda::ptx::prmt_b4e));));
+#endif // __cccl_ptx_isa >= 200
+
+#if __cccl_ptx_isa >= 200
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // prmt.b32.rc8 dest, a_reg, b_reg, c_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(
+                     static_cast<uint32_t (*)(int32_t, int32_t, uint32_t)>(cuda::ptx::prmt_rc8));));
+#endif // __cccl_ptx_isa >= 200
+
+#if __cccl_ptx_isa >= 200
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // prmt.b32.ecl dest, a_reg, b_reg, c_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(
+                     static_cast<uint32_t (*)(int32_t, int32_t, uint32_t)>(cuda::ptx::prmt_ecl));));
+#endif // __cccl_ptx_isa >= 200
+
+#if __cccl_ptx_isa >= 200
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // prmt.b32.ecr dest, a_reg, b_reg, c_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(
+                     static_cast<uint32_t (*)(int32_t, int32_t, uint32_t)>(cuda::ptx::prmt_ecr));));
+#endif // __cccl_ptx_isa >= 200
+
+#if __cccl_ptx_isa >= 200
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // prmt.b32.rc16 dest, a_reg, b_reg, c_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(
+                     static_cast<uint32_t (*)(int32_t, int32_t, uint32_t)>(cuda::ptx::prmt_rc16));));
+#endif // __cccl_ptx_isa >= 200
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/generated/shl.h b/libcudacxx/test/libcudacxx/cuda/ptx/generated/shl.h
new file mode 100644
index 00000000000..09ef491ed30
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/generated/shl.h
@@ -0,0 +1,39 @@
+// This file was automatically generated. Do not edit.
+
+// We use a special strategy to force the generation of the PTX. This is mainly
+// a fight against dead-code-elimination in the NVVM layer.
+//
+// The reason we need this strategy is because certain older versions of ptxas
+// segfault when a non-sensical sequence of PTX is generated. So instead, we try
+// to force the instantiation and compilation to PTX of all the overloads of the
+// PTX wrapping functions.
+//
+// We do this by writing a function pointer of each overload to the kernel
+// parameter `fn_ptr`.
+//
+// Because `fn_ptr` is possibly visible outside this translation unit, the
+// compiler must compile all the functions which are stored.
+
+__global__ void test_shl(void** fn_ptr)
+{
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shl.b16 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int16_t (*)(int16_t, uint32_t)>(cuda::ptx::shl));));
+#endif // __cccl_ptx_isa >= 100
+
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shl.b32 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int32_t (*)(int32_t, uint32_t)>(cuda::ptx::shl));));
+#endif // __cccl_ptx_isa >= 100
+
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shl.b64 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int64_t (*)(int64_t, uint32_t)>(cuda::ptx::shl));));
+#endif // __cccl_ptx_isa >= 100
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/generated/shr.h b/libcudacxx/test/libcudacxx/cuda/ptx/generated/shr.h
new file mode 100644
index 00000000000..2f56b9758cd
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/generated/shr.h
@@ -0,0 +1,60 @@
+// This file was automatically generated. Do not edit.
+
+// We use a special strategy to force the generation of the PTX. This is mainly
+// a fight against dead-code-elimination in the NVVM layer.
+//
+// The reason we need this strategy is because certain older versions of ptxas
+// segfault when a non-sensical sequence of PTX is generated. So instead, we try
+// to force the instantiation and compilation to PTX of all the overloads of the
+// PTX wrapping functions.
+//
+// We do this by writing a function pointer of each overload to the kernel
+// parameter `fn_ptr`.
+//
+// Because `fn_ptr` is possibly visible outside this translation unit, the
+// compiler must compile all the functions which are stored.
+
+__global__ void test_shr(void** fn_ptr)
+{
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shr.b16 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int16_t (*)(int16_t, uint32_t)>(cuda::ptx::shr));));
+#endif // __cccl_ptx_isa >= 100
+
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shr.b32 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int32_t (*)(int32_t, uint32_t)>(cuda::ptx::shr));));
+#endif // __cccl_ptx_isa >= 100
+
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shr.b64 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int64_t (*)(int64_t, uint32_t)>(cuda::ptx::shr));));
+#endif // __cccl_ptx_isa >= 100
+
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shr.s16 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int16_t (*)(int16_t, uint32_t)>(cuda::ptx::shr));));
+#endif // __cccl_ptx_isa >= 100
+
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shr.s32 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int32_t (*)(int32_t, uint32_t)>(cuda::ptx::shr));));
+#endif // __cccl_ptx_isa >= 100
+
+#if __cccl_ptx_isa >= 100
+  NV_IF_TARGET(NV_PROVIDES_SM_50,
+               (
+                   // shr.s64 dest, a_reg, b_reg;
+                   * fn_ptr++ = reinterpret_cast<void*>(static_cast<int64_t (*)(int64_t, uint32_t)>(cuda::ptx::shr));));
+#endif // __cccl_ptx_isa >= 100
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.bmsk.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.bmsk.compile.pass.cpp
new file mode 100644
index 00000000000..3d9eb18d3e1
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.bmsk.compile.pass.cpp
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: libcpp-has-no-threads
+
+// <cuda/ptx>
+
+#include <cuda/ptx>
+#include <cuda/std/utility>
+
+#include "generated/bmsk.h"
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.prmt.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.prmt.compile.pass.cpp
new file mode 100644
index 00000000000..68acbb13c71
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.prmt.compile.pass.cpp
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: libcpp-has-no-threads
+
+// <cuda/ptx>
+
+#include <cuda/ptx>
+#include <cuda/std/utility>
+
+#include "generated/prmt.h"
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.shl.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.shl.compile.pass.cpp
new file mode 100644
index 00000000000..85b51bbcd86
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.shl.compile.pass.cpp
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: libcpp-has-no-threads
+
+// <cuda/ptx>
+
+#include <cuda/ptx>
+#include <cuda/std/utility>
+
+#include "generated/shl.h"
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/ptx.shr.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.shr.compile.pass.cpp
new file mode 100644
index 00000000000..310b361129b
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/ptx/ptx.shr.compile.pass.cpp
@@ -0,0 +1,22 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: libcpp-has-no-threads
+
+// <cuda/ptx>
+
+#include <cuda/ptx>
+#include <cuda/std/utility>
+
+#include "generated/shr.h"
+
+int main(int, char**)
+{
+  return 0;
+}

From 29d14d43382f76082dd5973818803aa791ab5f4f Mon Sep 17 00:00:00 2001
From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com>
Date: Tue, 25 Feb 2025 16:46:14 -0600
Subject: [PATCH 12/17] Add test_reduce_api.py::test_reduce_struct_type_minmax
 (#3938)

* Add test_reduce_api.py::test_reduce_struct_type_minmax

This adds an example of using struct that holds minimum and maximum values
to compute smallest and largest values of an array using a single call
cub::DeviceReduce

It also demonstrates a use of TransformedIterator to map input array
values as v -> MinMax(v,v) to make reduction binary operator commutative,
to render data-parallel algorithm applicable. Since this is done in
registers, each element of input array is only accessed once.

* Refined docstring text

* Change reduce-minmax to compute mininum/maximum absolute values of an array

* Replace embedding_op to transform_op, add comment
---
 .../cuda/parallel/experimental/struct.py      | 10 ++--
 python/cuda_parallel/tests/test_reduce_api.py | 58 +++++++++++++++++++
 2 files changed, 63 insertions(+), 5 deletions(-)

diff --git a/python/cuda_parallel/cuda/parallel/experimental/struct.py b/python/cuda_parallel/cuda/parallel/experimental/struct.py
index 8338ddf84de..bd9cfb4feaa 100644
--- a/python/cuda_parallel/cuda/parallel/experimental/struct.py
+++ b/python/cuda_parallel/cuda/parallel/experimental/struct.py
@@ -38,15 +38,15 @@ def gpu_struct(this: type) -> Type[GpuStruct]:
 
     Example:
         The code snippet below shows how to use `gpu_struct` to define
-        a `Pixel` type (composed of `r`, `g` and `b` values),  and perform
-        a reduction on an array of `Pixel` objects to identify the one
-        with the largest `g` component:
+        a `MinMax` type (composed of `min_val`, `max_val` values), and perform
+        a reduction on an input array of floating point values to compute its
+        the smallest and the largest absolute values:
 
         .. literalinclude:: ../../python/cuda_parallel/tests/test_reduce_api.py
             :language: python
             :dedent:
-            :start-after: example-begin reduce-struct
-            :end-before: example-end reduce-struct
+            :start-after: example-begin reduce-minmax
+            :end-before: example-end reduce-minmax
     """
     # Implementation-wise, @gpu_struct creates and registers a
     # corresponding numba type to the given type, so that it can be
diff --git a/python/cuda_parallel/tests/test_reduce_api.py b/python/cuda_parallel/tests/test_reduce_api.py
index 9a36a0012fc..6087a263ee3 100644
--- a/python/cuda_parallel/tests/test_reduce_api.py
+++ b/python/cuda_parallel/tests/test_reduce_api.py
@@ -213,3 +213,61 @@ def max_g_value(x, y):
 
     np.testing.assert_equal(expected["g"], d_out.get()["g"])
     # example-end reduce-struct
+
+
+def test_reduce_struct_type_minmax():
+    # example-begin reduce-minmax
+    import cupy as cp
+    import numpy as np
+
+    import cuda.parallel.experimental.algorithms as algorithms
+    import cuda.parallel.experimental.iterators as iterators
+    from cuda.parallel.experimental.struct import gpu_struct
+
+    @gpu_struct
+    class MinMax:
+        min_val: np.float64
+        max_val: np.float64
+
+    def minmax_op(v1: MinMax, v2: MinMax):
+        c_min = min(v1.min_val, v2.min_val)
+        c_max = max(v1.max_val, v2.max_val)
+        return MinMax(c_min, c_max)
+
+    def transform_op(v):
+        av = abs(v)
+        return MinMax(av, av)
+
+    nelems = 4096
+
+    d_in = cp.random.randn(nelems)
+    # input values must be transformed to MinMax structures
+    # in-place to map computation to data-parallel reduction
+    # algorithm that requires commutative binary operation
+    # with both operands having the same type.
+    tr_it = iterators.TransformIterator(d_in, transform_op)
+
+    d_out = cp.empty(tuple(), dtype=MinMax.dtype)
+
+    # initial value set with identity elements of
+    # minimum and maximum operators
+    h_init = MinMax(np.inf, -np.inf)
+
+    # get algorithm object
+    cccl_sum = algorithms.reduce_into(tr_it, d_out, minmax_op, h_init)
+
+    # allocated needed temporary
+    tmp_sz = cccl_sum(None, tr_it, d_out, nelems, h_init)
+    tmp_storage = cp.empty(tmp_sz, dtype=cp.uint8)
+
+    # invoke the reduction algorithm
+    cccl_sum(tmp_storage, tr_it, d_out, nelems, h_init)
+
+    # display values computed on the device
+    actual = d_out.get()
+
+    h = np.abs(d_in.get())
+    expected = np.asarray([(h.min(), h.max())], dtype=MinMax.dtype)
+
+    assert actual == expected
+    # example-end reduce-minmax

From 10e8f25a84d175db3d4b7facad9b19c71c0c6a85 Mon Sep 17 00:00:00 2001
From: Federico Busato <50413820+fbusato@users.noreply.github.com>
Date: Tue, 25 Feb 2025 22:54:20 -0800
Subject: [PATCH 13/17] Add `cuda::std::aligned_accessor` (#3731)

Co-authored-by: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Co-authored-by: Mark Hoemmen <mhoemmen@users.noreply.github.com>
---
 .../standard_api/container_library/mdspan.rst |  12 ++-
 .../cuda/std/__mdspan/aligned_accessor.h      | 101 ++++++++++++++++++
 .../std/__memory/is_sufficiently_aligned.h    |  39 +++++++
 libcudacxx/include/cuda/std/mdspan            |   1 +
 .../aligned_accessor.alignment.fail.cpp       |  36 +++++++
 .../aligned_accessor.pass.cpp                 |  72 +++++++++++++
 6 files changed, 258 insertions(+), 3 deletions(-)
 create mode 100644 libcudacxx/include/cuda/std/__mdspan/aligned_accessor.h
 create mode 100644 libcudacxx/include/cuda/std/__memory/is_sufficiently_aligned.h
 create mode 100644 libcudacxx/test/libcudacxx/std/containers/views/mdspan/mdspan.aligned_accessor/aligned_accessor.alignment.fail.cpp
 create mode 100644 libcudacxx/test/libcudacxx/std/containers/views/mdspan/mdspan.aligned_accessor/aligned_accessor.pass.cpp

diff --git a/docs/libcudacxx/standard_api/container_library/mdspan.rst b/docs/libcudacxx/standard_api/container_library/mdspan.rst
index 0cebc85d1c8..990c3782292 100644
--- a/docs/libcudacxx/standard_api/container_library/mdspan.rst
+++ b/docs/libcudacxx/standard_api/container_library/mdspan.rst
@@ -3,12 +3,18 @@
 ``<cuda/std/mdspan>``
 ======================
 
+Provided functionalities
+------------------------
+
+-  All features of ``<mdspan>`` are made available in C++17 onwards
+-  C++26 ``std::dims`` is made available in C++17 onwards
+-  C++26 ``std::aligned_accessor`` is made available in C++17 onwards
+
 Extensions
 ----------
 
--  All features of ``<mdspan>`` are made available in C++14 onwards
--  The C++23 multidimensional ``operator[]`` is replaced with ``operator()``
--  C++26 ``std::dims`` is made available in C++14 onwards
+-  The C++23 multidimensional ``operator[]`` is replaced with ``operator()`` in previous C++ standards
+-  Detection of out-of-bounds accesses is available in debug mode
 
 Restrictions
 ------------
diff --git a/libcudacxx/include/cuda/std/__mdspan/aligned_accessor.h b/libcudacxx/include/cuda/std/__mdspan/aligned_accessor.h
new file mode 100644
index 00000000000..6c47b9c3ab2
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__mdspan/aligned_accessor.h
@@ -0,0 +1,101 @@
+//@HEADER
+// ************************************************************************
+//
+//                        Kokkos v. 4.0
+//       Copyright (2022) National Technology & Engineering
+//               Solutions of Sandia, LLC (NTESS).
+//
+// Under the terms of Contract DE-NA0003525 with NTESS,
+// the U.S. Government retains certain rights in this software.
+//
+// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
+// See https://kokkos.org/LICENSE for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+// ************************************************************************
+//@HEADER
+
+#ifndef _LIBCUDACXX___MDSPAN_ALIGNED_ACCESSOR_H
+#define _LIBCUDACXX___MDSPAN_ALIGNED_ACCESSOR_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/version>
+
+#if defined(__cccl_lib_mdspan)
+
+#  include <cuda/std/__bit/has_single_bit.h>
+#  include <cuda/std/__concepts/concept_macros.h>
+#  include <cuda/std/__mdspan/default_accessor.h>
+#  include <cuda/std/__memory/assume_aligned.h>
+#  include <cuda/std/__type_traits/is_abstract.h>
+#  include <cuda/std/__type_traits/is_array.h>
+#  include <cuda/std/__type_traits/is_convertible.h>
+#  include <cuda/std/__type_traits/is_object.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+template <class _ElementType, size_t _ByteAlignment>
+class aligned_accessor
+{
+public:
+  static constexpr auto byte_alignment = _ByteAlignment;
+
+  static_assert(_CUDA_VSTD::has_single_bit(byte_alignment), "byte_alignment must be a power of two.");
+
+  static_assert(byte_alignment >= alignof(_ElementType), "Insufficient byte alignment for _ElementType");
+
+  static_assert(_CCCL_TRAIT(is_object, _ElementType) && !_CCCL_TRAIT(is_abstract, _ElementType)
+                  && !_CCCL_TRAIT(is_array, _ElementType),
+                "_ElementType must be a complete object type that is neither an abstract class type nor an array "
+                "type.");
+
+  using offset_policy    = default_accessor<_ElementType>;
+  using element_type     = _ElementType;
+  using reference        = _ElementType&;
+  using data_handle_type = _ElementType*;
+
+  _CCCL_HIDE_FROM_ABI aligned_accessor() noexcept = default;
+
+  _CCCL_TEMPLATE(class _OtherElementType, size_t _OtherByteAlignment)
+  _CCCL_REQUIRES(_CCCL_TRAIT(is_convertible, _OtherElementType (*)[], element_type (*)[])
+                   _CCCL_AND((_OtherByteAlignment >= byte_alignment)))
+  _LIBCUDACXX_HIDE_FROM_ABI constexpr aligned_accessor(aligned_accessor<_OtherElementType, _OtherByteAlignment>) noexcept
+  {}
+
+  _CCCL_TEMPLATE(class _OtherElementType)
+  _CCCL_REQUIRES(_CCCL_TRAIT(is_convertible, _OtherElementType (*)[], element_type (*)[]))
+  _LIBCUDACXX_HIDE_FROM_ABI constexpr aligned_accessor(default_accessor<_OtherElementType>) noexcept {}
+
+  _CCCL_TEMPLATE(class _OtherElementType)
+  _CCCL_REQUIRES(_CCCL_TRAIT(is_convertible, _OtherElementType (*)[], element_type (*)[]))
+  _LIBCUDACXX_HIDE_FROM_ABI constexpr operator default_accessor<_OtherElementType>() const noexcept
+  {
+    return {};
+  }
+
+  _LIBCUDACXX_HIDE_FROM_ABI constexpr reference access(data_handle_type __p, size_t __i) const noexcept
+  {
+    return _CUDA_VSTD::assume_aligned<byte_alignment>(__p)[__i];
+  }
+
+  _LIBCUDACXX_HIDE_FROM_ABI constexpr typename offset_policy::data_handle_type
+  offset(data_handle_type __p, size_t __i) const noexcept
+  {
+    return _CUDA_VSTD::assume_aligned<byte_alignment>(__p) + __i;
+  }
+};
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // defined(__cccl_lib_mdspan)
+#endif // _LIBCUDACXX___MDSPAN_ALIGNED_ACCESSOR_H
diff --git a/libcudacxx/include/cuda/std/__memory/is_sufficiently_aligned.h b/libcudacxx/include/cuda/std/__memory/is_sufficiently_aligned.h
new file mode 100644
index 00000000000..da64fee6ed2
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__memory/is_sufficiently_aligned.h
@@ -0,0 +1,39 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBCUDACXX___MEMORY_IS_SUFFICIENTLY_ALIGNED_H
+#define _LIBCUDACXX___MEMORY_IS_SUFFICIENTLY_ALIGNED_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__bit/bit_cast.h>
+#include <cuda/std/cstddef> // size_t
+#include <cuda/std/cstdint> // uintptr_t
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+template <size_t _ByteAlignment, class _ElementType>
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool is_sufficiently_aligned(_ElementType* __ptr) noexcept
+{
+  return _CUDA_VSTD::bit_cast<uintptr_t>(__ptr) % _ByteAlignment == 0;
+}
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _LIBCUDACXX___MEMORY_IS_SUFFICIENTLY_ALIGNED_H
diff --git a/libcudacxx/include/cuda/std/mdspan b/libcudacxx/include/cuda/std/mdspan
index 78af32335ad..dd608fc8679 100644
--- a/libcudacxx/include/cuda/std/mdspan
+++ b/libcudacxx/include/cuda/std/mdspan
@@ -21,6 +21,7 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/std/__mdspan/aligned_accessor.h>
 #include <cuda/std/__mdspan/default_accessor.h>
 #include <cuda/std/__mdspan/dynamic_extent.h>
 #include <cuda/std/__mdspan/extents.h>
diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/mdspan.aligned_accessor/aligned_accessor.alignment.fail.cpp b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/mdspan.aligned_accessor/aligned_accessor.alignment.fail.cpp
new file mode 100644
index 00000000000..7f4910ab088
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/mdspan.aligned_accessor/aligned_accessor.alignment.fail.cpp
@@ -0,0 +1,36 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: msvc && c++17
+
+#include <cuda/std/mdspan>
+
+#include <test_macros.h>
+
+int main(int, char**)
+{
+  using T = int;
+  // conversion from smaller aligned accessor
+  cuda::std::aligned_accessor<T, sizeof(T)> aligned_x1{};
+  cuda::std::aligned_accessor<T, sizeof(T) * 2> aligned_x2{aligned_x1};
+  unused(aligned_x2);
+
+  // alignment to small
+  cuda::std::aligned_accessor<T, sizeof(T) / 2> aligned_half{};
+  unused(aligned_half);
+
+  // alignment non-power of 2
+  cuda::std::aligned_accessor<T, 6> aligned6{};
+  unused(aligned_half);
+
+  // non-convertible
+  cuda::std::aligned_accessor<T, sizeof(T)> aligned_int{};
+  cuda::std::aligned_accessor<float, sizeof(T)> aligned_float{aligned_int};
+  unused(aligned_float);
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/containers/views/mdspan/mdspan.aligned_accessor/aligned_accessor.pass.cpp b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/mdspan.aligned_accessor/aligned_accessor.pass.cpp
new file mode 100644
index 00000000000..72b5484c082
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/containers/views/mdspan/mdspan.aligned_accessor/aligned_accessor.pass.cpp
@@ -0,0 +1,72 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: msvc && c++17
+
+#include <cuda/std/cassert>
+#include <cuda/std/mdspan>
+#include <cuda/std/type_traits>
+
+#include <test_macros.h>
+
+template <class ElementType>
+__host__ __device__ constexpr void take_default_accessor_generic(cuda::std::default_accessor<ElementType>)
+{}
+
+template <class ElementType>
+cuda::std::enable_if_t<cuda::std::is_const_v<ElementType>>
+  __host__ __device__ constexpr take_default_accessor_generic_const(cuda::std::default_accessor<ElementType>)
+{}
+
+__host__ __device__ constexpr void take_default_accessor(cuda::std::default_accessor<int>) {}
+
+__host__ __device__ constexpr void take_default_accessor_const(cuda::std::default_accessor<const int>) {}
+
+__host__ __device__ constexpr bool test()
+{
+  using T = int;
+  using E = cuda::std::extents<size_t, 2>;
+  using L = cuda::std::layout_right;
+  using A = cuda::std::aligned_accessor<T, sizeof(T)>;
+  A aligned_non_const{};
+  cuda::std::aligned_accessor<const T, sizeof(T)> aligned_const{};
+  // basic tests
+  cuda::std::array<T, 2> d{42, 43};
+  cuda::std::mdspan<T, E, L, A> md(d.data(), 2);
+  assert(md(0) == 42);
+  assert(md(1) == 43);
+  // offset
+  assert(aligned_non_const.offset(d.data(), 0) == d.data());
+  assert(aligned_non_const.offset(d.data(), 2) == d.data() + 2);
+
+  // different alignment
+  cuda::std::aligned_accessor<T, sizeof(T) * 2> aligned_x2{};
+  cuda::std::aligned_accessor<T, sizeof(T)> aligned_x1{aligned_x2};
+  unused(aligned_x1);
+
+  // aligned accessor <-> default accessor
+  cuda::std::default_accessor<int> acc1{aligned_non_const};
+  cuda::std::default_accessor<int> acc2 = aligned_non_const;
+  unused(acc1);
+  unused(acc2);
+  take_default_accessor(aligned_non_const);
+
+  cuda::std::default_accessor<const T> acc3{aligned_const};
+  cuda::std::default_accessor<const T> acc4 = aligned_const;
+  unused(acc3);
+  unused(acc4);
+  take_default_accessor_const(aligned_const);
+  return true;
+}
+
+int main(int, char**)
+{
+  test();
+  static_assert(test());
+  return 0;
+}

From 497843a98d0ec855c598b6280a86a2ce81bb2af9 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Wed, 26 Feb 2025 10:36:25 +0100
Subject: [PATCH 14/17] Drop deprecated features from CUB util_ptx.cuh (#3935)

---
 cub/cub/util_ptx.cuh | 269 -------------------------------------------
 1 file changed, 269 deletions(-)

diff --git a/cub/cub/util_ptx.cuh b/cub/cub/util_ptx.cuh
index 206e815a761..d53a6cdc70a 100644
--- a/cub/cub/util_ptx.cuh
+++ b/cub/cub/util_ptx.cuh
@@ -52,28 +52,6 @@ CUB_NAMESPACE_BEGIN
  * Inlined PTX intrinsics
  ******************************************************************************/
 
-/**
- * \brief Shift-right then add.  Returns (\p x >> \p shift) + \p addend.
- */
-CCCL_DEPRECATED_BECAUSE("will be removed in the next major release")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int SHR_ADD(unsigned int x, unsigned int shift, unsigned int addend)
-{
-  unsigned int ret;
-  asm("vshr.u32.u32.u32.clamp.add %0, %1, %2, %3;" : "=r"(ret) : "r"(x), "r"(shift), "r"(addend));
-  return ret;
-}
-
-/**
- * \brief Shift-left then add.  Returns (\p x << \p shift) + \p addend.
- */
-CCCL_DEPRECATED_BECAUSE("will be removed in the next major release")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int SHL_ADD(unsigned int x, unsigned int shift, unsigned int addend)
-{
-  unsigned int ret;
-  asm("vshl.u32.u32.u32.clamp.add %0, %1, %2, %3;" : "=r"(ret) : "r"(x), "r"(shift), "r"(addend));
-  return ret;
-}
-
 #ifndef _CCCL_DOXYGEN_INVOKED // Do not document
 
 /**
@@ -124,135 +102,8 @@ _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int BFE(UnsignedBits source, unsigned in
   return BFE(source, bit_start, num_bits, detail::constant_v<int{sizeof(UnsignedBits)}>);
 }
 
-/**
- * \brief Bitfield insert.  Inserts the \p num_bits least significant bits of \p y into \p x at bit-offset \p bit_start.
- */
-CCCL_DEPRECATED_BECAUSE("will be removed in the next major release")
-_CCCL_DEVICE _CCCL_FORCEINLINE void
-BFI(unsigned int& ret, unsigned int x, unsigned int y, unsigned int bit_start, unsigned int num_bits)
-{
-  asm("bfi.b32 %0, %1, %2, %3, %4;" : "=r"(ret) : "r"(y), "r"(x), "r"(bit_start), "r"(num_bits));
-}
-
-/**
- * \brief Three-operand add.  Returns \p x + \p y + \p z.
- */
-CCCL_DEPRECATED_BECAUSE("will be removed in the next major release")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int IADD3(unsigned int x, unsigned int y, unsigned int z)
-{
-  asm("vadd.u32.u32.u32.add %0, %1, %2, %3;" : "=r"(x) : "r"(x), "r"(y), "r"(z));
-  return x;
-}
-
-/**
- * \brief Byte-permute. Pick four arbitrary bytes from two 32-bit registers, and reassemble them into a 32-bit
- * destination register.  For SM2.0 or later.
- *
- * \par
- * The bytes in the two source registers \p a and \p b are numbered from 0 to 7:
- * {\p b, \p a} = {{b7, b6, b5, b4}, {b3, b2, b1, b0}}. For each of the four bytes
- * {b3, b2, b1, b0} selected in the return value, a 4-bit selector is defined within
- * the four lower "nibbles" of \p index: {\p index } = {n7, n6, n5, n4, n3, n2, n1, n0}
- *
- * \par Snippet
- * The code snippet below illustrates byte-permute.
- * \par
- * \code
- * #include <cub/cub.cuh>
- *
- * __global__ void ExampleKernel(...)
- * {
- *     int a        = 0x03020100;
- *     int b        = 0x07060504;
- *     int index    = 0x00007531;
- *
- *     int selected = PRMT(a, b, index);    // 0x07050301
- *
- * \endcode
- *
- */
-CCCL_DEPRECATED_BECAUSE("will be removed in the next major release")
-_CCCL_DEVICE _CCCL_FORCEINLINE int PRMT(unsigned int a, unsigned int b, unsigned int index)
-{
-  int ret;
-  asm("prmt.b32 %0, %1, %2, %3;" : "=r"(ret) : "r"(a), "r"(b), "r"(index));
-  return ret;
-}
-
 #ifndef _CCCL_DOXYGEN_INVOKED // Do not document
 
-/**
- * Sync-threads barrier.
- */
-CCCL_DEPRECATED_BECAUSE("will be removed in the next major release")
-_CCCL_DEVICE _CCCL_FORCEINLINE void BAR(int count)
-{
-  asm volatile("bar.sync 1, %0;" : : "r"(count));
-}
-
-/**
- * CTA barrier
- */
-CCCL_DEPRECATED_BECAUSE("use __syncthreads() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE void CTA_SYNC()
-{
-  __syncthreads();
-}
-
-/**
- * CTA barrier with predicate
- */
-CCCL_DEPRECATED_BECAUSE("use __syncthreads_and() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE int CTA_SYNC_AND(int p)
-{
-  return __syncthreads_and(p);
-}
-
-/**
- * CTA barrier with predicate
- */
-CCCL_DEPRECATED_BECAUSE("use __syncthreads_or() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE int CTA_SYNC_OR(int p)
-{
-  return __syncthreads_or(p);
-}
-
-/**
- * Warp barrier
- */
-CCCL_DEPRECATED_BECAUSE("use __syncwarp() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE void WARP_SYNC(unsigned int member_mask)
-{
-  __syncwarp(member_mask);
-}
-
-/**
- * Warp any
- */
-CCCL_DEPRECATED_BECAUSE("use __any_sync() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE int WARP_ANY(int predicate, unsigned int member_mask)
-{
-  return __any_sync(member_mask, predicate);
-}
-
-/**
- * Warp any
- */
-CCCL_DEPRECATED_BECAUSE("use __all_sync() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE int WARP_ALL(int predicate, unsigned int member_mask)
-{
-  return __all_sync(member_mask, predicate);
-}
-
-/**
- * Warp ballot
- */
-CCCL_DEPRECATED_BECAUSE("use __ballot_sync() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE int WARP_BALLOT(int predicate, unsigned int member_mask)
-{
-  return __ballot_sync(member_mask, predicate);
-}
-
 /**
  * Warp synchronous shfl_up
  */
@@ -277,50 +128,6 @@ SHFL_DOWN_SYNC(unsigned int word, int src_offset, int flags, unsigned int member
   return word;
 }
 
-/**
- * Warp synchronous shfl_idx
- */
-CCCL_DEPRECATED_BECAUSE("use __shfl_sync() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int
-SHFL_IDX_SYNC(unsigned int word, int src_lane, int flags, unsigned int member_mask)
-{
-  asm volatile("shfl.sync.idx.b32 %0, %1, %2, %3, %4;"
-               : "=r"(word)
-               : "r"(word), "r"(src_lane), "r"(flags), "r"(member_mask));
-  return word;
-}
-
-/**
- * Warp synchronous shfl_idx
- */
-CCCL_DEPRECATED_BECAUSE("use __shfl_sync() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int SHFL_IDX_SYNC(unsigned int word, int src_lane, unsigned int member_mask)
-{
-  return __shfl_sync(member_mask, word, src_lane);
-}
-
-/**
- * Floating point multiply. (Mantissa LSB rounds towards zero.)
- */
-CCCL_DEPRECATED_BECAUSE("will be removed in the next major release")
-_CCCL_DEVICE _CCCL_FORCEINLINE float FMUL_RZ(float a, float b)
-{
-  float d;
-  asm("mul.rz.f32 %0, %1, %2;" : "=f"(d) : "f"(a), "f"(b));
-  return d;
-}
-
-/**
- * Floating point multiply-add. (Mantissa LSB rounds towards zero.)
- */
-CCCL_DEPRECATED_BECAUSE("will be removed in the next major release")
-_CCCL_DEVICE _CCCL_FORCEINLINE float FFMA_RZ(float a, float b, float c)
-{
-  float d;
-  asm("fma.rz.f32 %0, %1, %2, %3;" : "=f"(d) : "f"(a), "f"(b), "f"(c));
-  return d;
-}
-
 #endif // _CCCL_DOXYGEN_INVOKED
 
 /**
@@ -331,15 +138,6 @@ _CCCL_DEVICE _CCCL_FORCEINLINE void ThreadExit()
   asm volatile("exit;");
 }
 
-/**
- * \brief  Abort execution and generate an interrupt to the host CPU
- */
-CCCL_DEPRECATED_BECAUSE("use cuda::std::terminate() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE void ThreadTrap()
-{
-  asm volatile("trap;");
-}
-
 /**
  * \brief Returns the row-major linear thread identifier for a multidimensional thread block
  */
@@ -349,29 +147,6 @@ _CCCL_DEVICE _CCCL_FORCEINLINE int RowMajorTid(int block_dim_x, int block_dim_y,
        + ((block_dim_y == 1) ? 0 : (threadIdx.y * block_dim_x)) + threadIdx.x;
 }
 
-/**
- * \brief Returns the warp lane ID of the calling thread
- */
-CCCL_DEPRECATED_BECAUSE("use cuda::ptx::get_sreg_laneid() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int LaneId()
-{
-  unsigned int ret;
-  asm("mov.u32 %0, %%laneid;" : "=r"(ret));
-  return ret;
-}
-
-/**
- * \brief Returns the warp ID of the calling thread.  Warp ID is guaranteed to be unique among warps, but may not
- * correspond to a zero-based ranking within the thread block.
- */
-CCCL_DEPRECATED_BECAUSE("use cuda::ptx::get_sreg_warpid() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int WarpId()
-{
-  unsigned int ret;
-  asm("mov.u32 %0, %%warpid;" : "=r"(ret));
-  return ret;
-}
-
 /**
  * @brief Returns the warp mask for a warp of @p LOGICAL_WARP_THREADS threads
  *
@@ -401,50 +176,6 @@ _CCCL_HOST_DEVICE _CCCL_FORCEINLINE unsigned int WarpMask(unsigned int warp_id)
   return member_mask;
 }
 
-/**
- * \brief Returns the warp lane mask of all lanes less than the calling thread
- */
-CCCL_DEPRECATED_BECAUSE("use cuda::ptx::get_sreg_lanemask_lt() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int LaneMaskLt()
-{
-  unsigned int ret;
-  asm("mov.u32 %0, %%lanemask_lt;" : "=r"(ret));
-  return ret;
-}
-
-/**
- * \brief Returns the warp lane mask of all lanes less than or equal to the calling thread
- */
-CCCL_DEPRECATED_BECAUSE("use cuda::ptx::get_sreg_lanemask_le() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int LaneMaskLe()
-{
-  unsigned int ret;
-  asm("mov.u32 %0, %%lanemask_le;" : "=r"(ret));
-  return ret;
-}
-
-/**
- * \brief Returns the warp lane mask of all lanes greater than the calling thread
- */
-CCCL_DEPRECATED_BECAUSE("use cuda::ptx::get_sreg_lanemask_gt() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int LaneMaskGt()
-{
-  unsigned int ret;
-  asm("mov.u32 %0, %%lanemask_gt;" : "=r"(ret));
-  return ret;
-}
-
-/**
- * \brief Returns the warp lane mask of all lanes greater than or equal to the calling thread
- */
-CCCL_DEPRECATED_BECAUSE("use cuda::ptx::get_sreg_lanemask_ge() instead")
-_CCCL_DEVICE _CCCL_FORCEINLINE unsigned int LaneMaskGe()
-{
-  unsigned int ret;
-  asm("mov.u32 %0, %%lanemask_ge;" : "=r"(ret));
-  return ret;
-}
-
 /**
  * @brief Shuffle-up for any data type.
  *        Each <em>warp-lane<sub>i</sub></em> obtains the value @p input contributed by

From 7fbbd247b85d97bda0cb224c3da95004a8de068c Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?C=C3=A9dric=20Augonnet?=
 <158148890+caugonnet@users.noreply.github.com>
Date: Wed, 26 Feb 2025 12:52:20 +0100
Subject: [PATCH 15/17] [STF] Thread safe graph_ctx (#3925)

* Introduce a mutex to protect the underlying CUDA graph of a graph context so that we can generate tasks concurrently

* Protect CUDA graphs against concurrent accesses

* do test results

* remove dead code

* Add a test with graph capture and threads

* Add and use with_locked_graph, also use weak_ptr

* Big code simplification ! no need for shared_ptr/weak_ptr with the mutex that will outlive its users

* use std::reference_wrapper to have graph_task be moved assignable

* replace complicated API based on lambda function with a simple lock guard

* restore a comment removed by mistake

* capture with the lock taken because we need to ensure the captured stream is not used concurrently

* Fix build

* comment why we put a reference_wrapper

* Add a sanity check to ensure we have finished all tasks

* atomic variables are not automatically initialized, so we do set them

* Add missing mutex headers

* improve readability

---------

Co-authored-by: Andrei Alexandrescu <andrei@erdani.com>
---
 .../experimental/__stf/graph/graph_ctx.cuh    |  8 +-
 .../experimental/__stf/graph/graph_task.cuh   | 44 +++++++++-
 .../__stf/internal/acquire_release.cuh        |  4 +
 .../__stf/internal/backend_ctx.cuh            | 30 +++++--
 .../__stf/internal/parallel_for_scope.cuh     | 14 +--
 cudax/test/stf/CMakeLists.txt                 |  2 +
 .../stf/threads/axpy-threads-graph-capture.cu | 85 +++++++++++++++++++
 cudax/test/stf/threads/axpy-threads-graph.cu  | 74 ++++++++++++++++
 8 files changed, 244 insertions(+), 17 deletions(-)
 create mode 100644 cudax/test/stf/threads/axpy-threads-graph-capture.cu
 create mode 100644 cudax/test/stf/threads/axpy-threads-graph.cu

diff --git a/cudax/include/cuda/experimental/__stf/graph/graph_ctx.cuh b/cudax/include/cuda/experimental/__stf/graph/graph_ctx.cuh
index ee76bdf8793..58ca20b0873 100644
--- a/cudax/include/cuda/experimental/__stf/graph/graph_ctx.cuh
+++ b/cudax/include/cuda/experimental/__stf/graph/graph_ctx.cuh
@@ -34,6 +34,8 @@
 #include <cuda/experimental/__stf/internal/parallel_for_scope.cuh>
 #include <cuda/experimental/__stf/places/blocked_partition.cuh> // for unit test!
 
+#include <mutex>
+
 namespace cuda::experimental::stf
 {
 
@@ -261,6 +263,9 @@ class graph_ctx : public backend_ctx<graph_ctx>
     bool submitted                                = false; // did we submit ?
     mutable bool explicit_graph                   = false;
 
+    // To protect _graph against concurrent modifications
+    ::std::mutex graph_mutex;
+
     executable_graph_cache_stat cache_stats;
 
     /* By default, the finalize operation is blocking, unless user provided
@@ -311,7 +316,8 @@ public:
   auto task(exec_place e_place, task_dep<Deps>... deps)
   {
     auto dump_hooks = reserved::get_dump_hooks(this, deps...);
-    auto result     = graph_task<Deps...>(*this, get_graph(), get_graph_epoch(), mv(e_place), mv(deps)...);
+    auto result =
+      graph_task<Deps...>(*this, get_graph(), this->state().graph_mutex, get_graph_epoch(), mv(e_place), mv(deps)...);
     result.add_post_submission_hook(dump_hooks);
     return result;
   }
diff --git a/cudax/include/cuda/experimental/__stf/graph/graph_task.cuh b/cudax/include/cuda/experimental/__stf/graph/graph_task.cuh
index 4066315bded..42b0bdf8e66 100644
--- a/cudax/include/cuda/experimental/__stf/graph/graph_task.cuh
+++ b/cudax/include/cuda/experimental/__stf/graph/graph_task.cuh
@@ -33,6 +33,8 @@
 #include <cuda/experimental/__stf/internal/logical_data.cuh>
 #include <cuda/experimental/__stf/internal/void_interface.cuh>
 
+#include <mutex>
+
 namespace cuda::experimental::stf
 {
 
@@ -57,9 +59,14 @@ public:
   // A cudaGraph_t is needed
   graph_task() = delete;
 
-  graph_task(backend_ctx_untyped ctx, cudaGraph_t g, size_t epoch, exec_place e_place = exec_place::current_device())
+  graph_task(backend_ctx_untyped ctx,
+             cudaGraph_t g,
+             ::std::mutex& graph_mutex,
+             size_t epoch,
+             exec_place e_place = exec_place::current_device())
       : task(mv(e_place))
       , ctx_graph(EXPECT(g))
+      , graph_mutex(graph_mutex)
       , epoch(epoch)
       , ctx(mv(ctx))
   {
@@ -75,6 +82,8 @@ public:
 
   graph_task& start()
   {
+    ::std::lock_guard<::std::mutex> lock(graph_mutex);
+
     event_list prereqs = acquire(ctx);
 
     // The CUDA graph API does not like duplicate dependencies
@@ -98,6 +107,8 @@ public:
   /* End the task, but do not clear its data structures yet */
   graph_task<>& end_uncleared()
   {
+    ::std::lock_guard<::std::mutex> lock(graph_mutex);
+
     cudaGraphNode_t n;
 
     auto done_prereqs = event_list();
@@ -388,6 +399,11 @@ public:
     return ctx_graph;
   }
 
+  [[nodiscard]] auto lock_ctx_graph()
+  {
+    return ::std::unique_lock<::std::mutex>(graph_mutex);
+  }
+
   void set_current_place(pos4 p)
   {
     get_exec_place().as_grid().set_current_place(ctx, p);
@@ -428,7 +444,15 @@ private:
 
   /* This is the support graph associated to the entire context */
   cudaGraph_t ctx_graph = nullptr;
-  size_t epoch          = 0;
+
+  // This protects ctx_graph : it's ok to store a reference to it because the
+  // context and this mutex will outlive the moment when this mutex is needed
+  // (and most likely the graph_task object)
+  // Note that we use a reference_wrapper instead of a mere reference to ensure
+  // the graph_task class remains move assignable.
+  ::std::reference_wrapper<::std::mutex> graph_mutex;
+
+  size_t epoch = 0;
 
   ::std::vector<cudaGraphNode_t> ready_dependencies;
 
@@ -448,8 +472,13 @@ template <typename... Deps>
 class graph_task : public graph_task<>
 {
 public:
-  graph_task(backend_ctx_untyped ctx, cudaGraph_t g, size_t epoch, exec_place e_place, task_dep<Deps>... deps)
-      : graph_task<>(mv(ctx), g, epoch, mv(e_place))
+  graph_task(backend_ctx_untyped ctx,
+             cudaGraph_t g,
+             ::std::mutex& graph_mutex,
+             size_t epoch,
+             exec_place e_place,
+             task_dep<Deps>... deps)
+      : graph_task<>(mv(ctx), g, graph_mutex, epoch, mv(e_place))
   {
     static_assert(sizeof(*this) == sizeof(graph_task<>), "Cannot add state - it would be lost by slicing.");
     add_deps(mv(deps)...);
@@ -544,6 +573,13 @@ public:
       // CAPTURE the lambda
       //
 
+      // To ensure the same CUDA stream is not used in multiple threads, we
+      // ensure there can't be multiple threads capturing at the same time.
+      //
+      // TODO : provide a per-thread CUDA stream dedicated for capture on that
+      // execution place.
+      auto lock = lock_ctx_graph();
+
       // Get a stream from the pool associated to the execution place
       cudaStream_t capture_stream = get_exec_place().getStream(ctx.async_resources(), true).stream;
 
diff --git a/cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh b/cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh
index 1ed35dceebc..d42286ea960 100644
--- a/cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh
+++ b/cudax/include/cuda/experimental/__stf/internal/acquire_release.cuh
@@ -324,6 +324,10 @@ inline void task::release(backend_ctx_untyped& ctx, event_list& done_prereqs)
   // This will, in particular, release shared_ptr to logical data captured in
   // the dependencies.
   pimpl->post_submission_hooks.clear();
+
+#ifndef NDEBUG
+  ctx.increment_finished_task_count();
+#endif
 }
 
 } // namespace cuda::experimental::stf
diff --git a/cudax/include/cuda/experimental/__stf/internal/backend_ctx.cuh b/cudax/include/cuda/experimental/__stf/internal/backend_ctx.cuh
index 79f4dfc3cac..56f4a719a4b 100644
--- a/cudax/include/cuda/experimental/__stf/internal/backend_ctx.cuh
+++ b/cudax/include/cuda/experimental/__stf/internal/backend_ctx.cuh
@@ -217,7 +217,9 @@ public:
     if constexpr (::std::is_same_v<Ctx, graph_ctx>)
     {
       cudaHostNodeParams params = {.fn = callback, .userData = wrapper};
+
       // Put this host node into the child graph that implements the graph_task<>
+      auto lock = t.lock_ctx_graph();
       cuda_safe_call(cudaGraphAddHostNode(&t.get_node(), t.get_ctx_graph(), nullptr, 0, &params));
     }
     else
@@ -410,19 +412,20 @@ public:
 
       if constexpr (::std::is_same_v<Ctx, graph_ctx>)
       {
+        auto lock = t.lock_ctx_graph();
+        auto& g   = t.get_ctx_graph();
+
         // We have two situations : either there is a single kernel and we put the kernel in the context's
         // graph, or we rely on a child graph
         if (res.size() == 1)
         {
-          insert_one_kernel(res[0], t.get_node(), t.get_ctx_graph());
+          insert_one_kernel(res[0], t.get_node(), g);
         }
         else
         {
           ::std::vector<cudaGraphNode_t>& chain = t.get_node_chain();
           chain.resize(res.size());
 
-          auto& g = t.get_ctx_graph();
-
           // Create a chain of kernels
           for (size_t i = 0; i < res.size(); i++)
           {
@@ -455,6 +458,7 @@ public:
 
       if constexpr (::std::is_same_v<Ctx, graph_ctx>)
       {
+        auto lock = t.lock_ctx_graph();
         insert_one_kernel(res, t.get_node(), t.get_ctx_graph());
       }
       else
@@ -561,8 +565,9 @@ protected:
 
     virtual ~impl()
     {
-      // We can't assert here because there may be tasks inside tasks
-      //_CCCL_ASSERT(total_task_cnt == 0, "You created some tasks but forgot to call finalize().");
+#ifndef NDEBUG
+      _CCCL_ASSERT(total_task_cnt == total_finished_task_cnt, "Not all tasks were finished.");
+#endif
 
       if (!is_recording_stats)
       {
@@ -708,7 +713,6 @@ protected:
     {
       // assert(!stack.hasCurrentTask());
       attached_allocators.clear();
-      total_task_cnt.store(0);
       // Leave custom_allocator, auto_scheduler, and auto_reordered as they were.
     }
 
@@ -732,7 +736,12 @@ protected:
       transfers;
     bool is_recording_stats = false;
     // Keep track of the number of tasks generated in the context
-    ::std::atomic<size_t> total_task_cnt;
+    ::std::atomic<size_t> total_task_cnt = 0;
+
+#ifndef NDEBUG
+    // Keep track of the number of completed tasks in that context
+    ::std::atomic<size_t> total_finished_task_cnt = 0;
+#endif
 
     // This data structure contains all resources useful for an efficient
     // asynchronous execution. This will for example contain pools of CUDA
@@ -851,6 +860,13 @@ public:
     ++pimpl->total_task_cnt;
   }
 
+#ifndef NDEBUG
+  void increment_finished_task_count()
+  {
+    ++pimpl->total_finished_task_cnt;
+  }
+#endif
+
   size_t task_count() const
   {
     return pimpl->total_task_cnt;
diff --git a/cudax/include/cuda/experimental/__stf/internal/parallel_for_scope.cuh b/cudax/include/cuda/experimental/__stf/internal/parallel_for_scope.cuh
index 06f1f7f689b..7016a35f9fd 100644
--- a/cudax/include/cuda/experimental/__stf/internal/parallel_for_scope.cuh
+++ b/cudax/include/cuda/experimental/__stf/internal/parallel_for_scope.cuh
@@ -698,7 +698,8 @@ public:
         kernel_params.sharedMemBytes = 0;
 
         // This new node will depend on the previous in the chain (allocation)
-        cuda_safe_call(cudaGraphAddKernelNode(&t.get_node(), t.get_ctx_graph(), NULL, 0, &kernel_params));
+        auto lock = t.lock_ctx_graph();
+        cudaGraphAddKernelNode(&t.get_node(), t.get_ctx_graph(), NULL, 0, &kernel_params);
       }
 
       return;
@@ -715,7 +716,8 @@ public:
       return ::std::pair(size_t(minGridSize), size_t(blockSize));
     }();
 
-    const auto [block_size, min_blocks] = conf;
+    const auto block_size = conf.first;
+    const auto min_blocks = conf.second;
 
     // max_blocks is computed so we have one thread per element processed
     const auto max_blocks = (n + block_size - 1) / block_size;
@@ -759,10 +761,8 @@ public:
     }
     else
     {
-      auto g = t.get_ctx_graph();
-
       _CCCL_ASSERT(sub_exec_place.is_device(), "Invalid execution place");
-      int dev_id = device_ordinal(sub_exec_place.affine_data_place());
+      const int dev_id = device_ordinal(sub_exec_place.affine_data_place());
 
       cudaMemAllocNodeParams allocParams{};
       allocParams.poolProps.allocType   = cudaMemAllocationTypePinned;
@@ -770,6 +770,8 @@ public:
       allocParams.poolProps.location    = {.type = cudaMemLocationTypeDevice, .id = dev_id};
       allocParams.bytesize              = blocks * sizeof(redux_vars<deps_tup_t, ops_and_inits>);
 
+      auto lock               = t.lock_ctx_graph();
+      auto g                  = t.get_ctx_graph();
       const auto& input_nodes = t.get_ready_dependencies();
 
       /* This first node depends on task's dependencies themselves */
@@ -898,7 +900,9 @@ public:
       // This task corresponds to a single graph node, so we set that
       // node instead of creating an child graph. Input and output
       // dependencies will be filled later.
+      auto lock = t.lock_ctx_graph();
       cuda_safe_call(cudaGraphAddKernelNode(&t.get_node(), t.get_ctx_graph(), nullptr, 0, &kernel_params));
+
       // fprintf(stderr, "KERNEL NODE => graph %p, gridDim %d blockDim %d (n %ld)\n", t.get_graph(),
       // kernel_params.gridDim.x, kernel_params.blockDim.x, n);
     }
diff --git a/cudax/test/stf/CMakeLists.txt b/cudax/test/stf/CMakeLists.txt
index b80c11f1cf1..f8765f19bbc 100644
--- a/cudax/test/stf/CMakeLists.txt
+++ b/cudax/test/stf/CMakeLists.txt
@@ -139,6 +139,8 @@ set(stf_test_codegen_sources
   stress/launch_vs_parallelfor.cu
   stress/parallel_for_overhead.cu
   # threads/axpy-threads-pfor.cu   # Currently has a difficult-to-reproduce concurrency problem
+  threads/axpy-threads-graph.cu
+  threads/axpy-threads-graph-capture.cu
   tools/auto_dump/auto_dump.cu
 )
 
diff --git a/cudax/test/stf/threads/axpy-threads-graph-capture.cu b/cudax/test/stf/threads/axpy-threads-graph-capture.cu
new file mode 100644
index 00000000000..6eabe56bc0b
--- /dev/null
+++ b/cudax/test/stf/threads/axpy-threads-graph-capture.cu
@@ -0,0 +1,85 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDASTF in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+/**
+ * @file
+ *
+ * @brief Ensure a graph_ctx can be used concurrently
+ *
+ */
+
+#include <cuda/experimental/stf.cuh>
+
+#include <mutex>
+#include <thread>
+
+using namespace cuda::experimental::stf;
+
+__global__ void axpy(slice<int> y, slice<const int> x, int a)
+{
+  int tid      = blockIdx.x * blockDim.x + threadIdx.x;
+  int nthreads = gridDim.x * blockDim.x;
+
+  for (int i = tid; i < x.size(); i += nthreads)
+  {
+    y(i) += a * x(i);
+  }
+}
+
+void mytask(graph_ctx ctx, int /*id*/)
+{
+  const size_t N = 16;
+
+  int alpha = 3;
+
+  auto lX = ctx.logical_data<int>(N);
+  auto lY = ctx.logical_data<int>(N);
+
+  ctx.parallel_for(lX.shape(), lX.write(), lY.write())->*[] __device__(size_t i, auto x, auto y) {
+    x(i) = (1 + i);
+    y(i) = (2 + i * i);
+  };
+
+  /* Compute Y = Y + alpha X */
+  for (size_t i = 0; i < 200; i++)
+  {
+    ctx.task(lY.rw(), lX.read())->*[alpha](cudaStream_t stream, auto dY, auto dX) {
+      axpy<<<128, 64, 0, stream>>>(dY, dX, alpha);
+    };
+  }
+
+  ctx.host_launch(lX.read(), lY.read())->*[alpha](auto x, auto y) {
+    for (size_t i = 0; i < N; i++)
+    {
+      EXPECT(x(i) == 1 + i);
+      EXPECT(y(i) == 2 + i * i + 200 * alpha * x(i));
+    }
+  };
+}
+
+int main()
+{
+  graph_ctx ctx;
+
+  ::std::vector<::std::thread> threads;
+  // Launch threads
+  for (int i = 0; i < 10; ++i)
+  {
+    threads.emplace_back(mytask, ctx, i);
+  }
+
+  // Wait for all threads to complete.
+  for (auto& th : threads)
+  {
+    th.join();
+  }
+
+  ctx.finalize();
+}
diff --git a/cudax/test/stf/threads/axpy-threads-graph.cu b/cudax/test/stf/threads/axpy-threads-graph.cu
new file mode 100644
index 00000000000..79ab5e91ce7
--- /dev/null
+++ b/cudax/test/stf/threads/axpy-threads-graph.cu
@@ -0,0 +1,74 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDASTF in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+/**
+ * @file
+ *
+ * @brief Ensure a graph_ctx can be used concurrently
+ *
+ */
+
+#include <cuda/experimental/stf.cuh>
+
+#include <mutex>
+#include <thread>
+
+using namespace cuda::experimental::stf;
+
+void mytask(graph_ctx ctx, int /*id*/)
+{
+  const size_t N = 16;
+
+  int alpha = 3;
+
+  auto lX = ctx.logical_data<int>(N);
+  auto lY = ctx.logical_data<int>(N);
+
+  ctx.parallel_for(lX.shape(), lX.write(), lY.write())->*[] __device__(size_t i, auto x, auto y) {
+    x(i) = (1 + i);
+    y(i) = (2 + i * i);
+  };
+
+  /* Compute Y = Y + alpha X */
+  for (size_t i = 0; i < 200; i++)
+  {
+    ctx.parallel_for(lY.shape(), lY.rw(), lX.read())->*[alpha] __device__(size_t i, auto dY, auto dX) {
+      dY(i) += alpha * dX(i);
+    };
+  }
+
+  ctx.host_launch(lX.read(), lY.read())->*[alpha](auto x, auto y) {
+    for (size_t i = 0; i < N; i++)
+    {
+      EXPECT(x(i) == 1 + i);
+      EXPECT(y(i) == 2 + i * i + 200 * alpha * x(i));
+    }
+  };
+}
+
+int main()
+{
+  graph_ctx ctx;
+
+  ::std::vector<::std::thread> threads;
+  // Launch threads
+  for (int i = 0; i < 10; ++i)
+  {
+    threads.emplace_back(mytask, ctx, i);
+  }
+
+  // Wait for all threads to complete.
+  for (auto& th : threads)
+  {
+    th.join();
+  }
+
+  ctx.finalize();
+}

From a52e6263d39bee5828c48a6c9d1de484c60bc382 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Wed, 26 Feb 2025 15:17:59 +0100
Subject: [PATCH 16/17] Replace CUB macros in tunings and benchmarks (#3931)

---
 c/parallel/src/reduce.cu                      |  5 +--
 cub/benchmarks/bench/partition/flagged.cu     |  3 +-
 cub/benchmarks/bench/partition/if.cu          |  3 +-
 cub/benchmarks/bench/select/flagged.cu        |  4 ++-
 cub/benchmarks/bench/select/if.cu             |  4 ++-
 cub/benchmarks/bench/select/unique.cu         |  6 ++--
 .../tuning/tuning_run_length_encode.cuh       | 10 +++---
 .../dispatch/tuning/tuning_select_if.cuh      |  5 +--
 thrust/testing/zip_iterator.cu                | 32 +++++++++----------
 9 files changed, 41 insertions(+), 31 deletions(-)

diff --git a/c/parallel/src/reduce.cu b/c/parallel/src/reduce.cu
index 383cfc0a895..4b458022c33 100644
--- a/c/parallel/src/reduce.cu
+++ b/c/parallel/src/reduce.cu
@@ -14,6 +14,7 @@
 #include <cub/grid/grid_even_share.cuh>
 #include <cub/util_device.cuh>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/cstdint>
 #include <cuda/std/functional> // ::cuda::std::identity
 #include <cuda/std/variant>
@@ -97,8 +98,8 @@ reduce_runtime_tuning_policy get_policy(int cc, cccl_type_info accumulator_type)
   auto [_, block_size, items_per_thread, vector_load_length] = find_tuning(cc, chain);
 
   // Implement part of MemBoundScaling
-  items_per_thread = CUB_MAX(1, CUB_MIN(items_per_thread * 4 / accumulator_type.size, items_per_thread * 2));
-  block_size       = CUB_MIN(block_size, (((1024 * 48) / (accumulator_type.size * items_per_thread)) + 31) / 32 * 32);
+  items_per_thread = cuda::std::clamp(items_per_thread * 4 / accumulator_type.size, 1, items_per_thread * 2);
+  block_size = _CUDA_VSTD::min(block_size, (((1024 * 48) / (accumulator_type.size * items_per_thread)) + 31) / 32 * 32);
 
   return {block_size, items_per_thread, vector_load_length};
 }
diff --git a/cub/benchmarks/bench/partition/flagged.cu b/cub/benchmarks/bench/partition/flagged.cu
index 7217ee32e6e..0a41f88f1a4 100644
--- a/cub/benchmarks/bench/partition/flagged.cu
+++ b/cub/benchmarks/bench/partition/flagged.cu
@@ -29,6 +29,7 @@
 
 #include <thrust/count.h>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/type_traits>
 
 #include <look_back_helper.cuh>
@@ -63,7 +64,7 @@ struct policy_hub_t
     static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = TUNE_ITEMS_PER_THREAD;
 
     static constexpr int ITEMS_PER_THREAD =
-      CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT))));
+      _CUDA_VSTD::clamp(NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT), 1, NOMINAL_4B_ITEMS_PER_THREAD);
 
     using SelectIfPolicyT =
       cub::AgentSelectIfPolicy<TUNE_THREADS_PER_BLOCK,
diff --git a/cub/benchmarks/bench/partition/if.cu b/cub/benchmarks/bench/partition/if.cu
index b5ed0b41b1a..5dc240fb2e7 100644
--- a/cub/benchmarks/bench/partition/if.cu
+++ b/cub/benchmarks/bench/partition/if.cu
@@ -29,6 +29,7 @@
 
 #include <thrust/count.h>
 
+#include <cuda/std/__algorithm_>
 #include <cuda/std/type_traits>
 
 #include <look_back_helper.cuh>
@@ -63,7 +64,7 @@ struct policy_hub_t
     static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = TUNE_ITEMS_PER_THREAD;
 
     static constexpr int ITEMS_PER_THREAD =
-      CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT))));
+      _CUDA_VSTD::clamp(NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT), 1, NOMINAL_4B_ITEMS_PER_THREAD);
 
     using SelectIfPolicyT =
       cub::AgentSelectIfPolicy<TUNE_THREADS_PER_BLOCK,
diff --git a/cub/benchmarks/bench/select/flagged.cu b/cub/benchmarks/bench/select/flagged.cu
index 2562cec6c76..168e2bc0cc5 100644
--- a/cub/benchmarks/bench/select/flagged.cu
+++ b/cub/benchmarks/bench/select/flagged.cu
@@ -29,6 +29,8 @@
 
 #include <thrust/count.h>
 
+#include <cuda/std/__algorithm_>
+
 #include <look_back_helper.cuh>
 #include <nvbench_helper.cuh>
 
@@ -61,7 +63,7 @@ struct policy_hub_t
     static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = TUNE_ITEMS_PER_THREAD;
 
     static constexpr int ITEMS_PER_THREAD =
-      CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT))));
+      _CUDA_VSTD::clamp(NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT), 1, NOMINAL_4B_ITEMS_PER_THREAD);
 
     using SelectIfPolicyT =
       cub::AgentSelectIfPolicy<TUNE_THREADS_PER_BLOCK,
diff --git a/cub/benchmarks/bench/select/if.cu b/cub/benchmarks/bench/select/if.cu
index a36cf256434..98bdec915a7 100644
--- a/cub/benchmarks/bench/select/if.cu
+++ b/cub/benchmarks/bench/select/if.cu
@@ -29,6 +29,8 @@
 
 #include <thrust/count.h>
 
+#include <cuda/std/__algorithm_>
+
 #include <limits>
 
 #include <look_back_helper.cuh>
@@ -63,7 +65,7 @@ struct policy_hub_t
     static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = TUNE_ITEMS_PER_THREAD;
 
     static constexpr int ITEMS_PER_THREAD =
-      CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT))));
+      _CUDA_VSTD::clamp(NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT), 1, NOMINAL_4B_ITEMS_PER_THREAD);
 
     using SelectIfPolicyT =
       cub::AgentSelectIfPolicy<TUNE_THREADS_PER_BLOCK,
diff --git a/cub/benchmarks/bench/select/unique.cu b/cub/benchmarks/bench/select/unique.cu
index 88f469ae953..75c0c3771e0 100644
--- a/cub/benchmarks/bench/select/unique.cu
+++ b/cub/benchmarks/bench/select/unique.cu
@@ -3,6 +3,8 @@
 
 #include <cub/device/device_select.cuh>
 
+#include <cuda/std/__algorithm_>
+
 #include <limits>
 
 #include <look_back_helper.cuh>
@@ -36,8 +38,8 @@ struct policy_hub_t
   {
     static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = TUNE_ITEMS_PER_THREAD;
 
-    static constexpr int ITEMS_PER_THREAD =
-      CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT))));
+    static constexpr int ITEMS_PER_THREAD = _CUDA_VSTD::min(
+      NOMINAL_4B_ITEMS_PER_THREAD, _CUDA_VSTD::max(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT))));
 
     using SelectIfPolicyT =
       cub::AgentSelectIfPolicy<TUNE_THREADS_PER_BLOCK,
diff --git a/cub/cub/device/dispatch/tuning/tuning_run_length_encode.cuh b/cub/cub/device/dispatch/tuning/tuning_run_length_encode.cuh
index 0e46a47d49d..9f272961ef2 100644
--- a/cub/cub/device/dispatch/tuning/tuning_run_length_encode.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_run_length_encode.cuh
@@ -46,7 +46,8 @@
 #include <cub/util_device.cuh>
 #include <cub/util_type.cuh>
 
-#include <cuda/std/__algorithm/max.h>
+#include <cuda/cmath>
+#include <cuda/std/__algorithm_>
 
 CUB_NAMESPACE_BEGIN
 
@@ -315,9 +316,8 @@ struct policy_hub
     static constexpr int items =
       (max_input_bytes <= 8)
         ? 6
-        // TODO(bgruber): use clamp() and ceil_div in C++14
-        : CUB_MIN(nominal_4B_items_per_thread,
-                  CUB_MAX(1, ((nominal_4B_items_per_thread * 8) + combined_input_bytes - 1) / combined_input_bytes));
+        : ::cuda::std::clamp(
+            ::cuda::ceil_div(nominal_4B_items_per_thread * 8, combined_input_bytes), 1, nominal_4B_items_per_thread);
     using ReduceByKeyPolicyT =
       AgentReduceByKeyPolicy<128,
                              items,
@@ -603,7 +603,7 @@ struct policy_hub
     static constexpr int nominal_4B_items_per_thread = 15;
     // TODO(bgruber): use clamp() in C++14
     static constexpr int ITEMS_PER_THREAD =
-      CUB_MIN(nominal_4B_items_per_thread, CUB_MAX(1, (nominal_4B_items_per_thread * 4 / sizeof(KeyT))));
+      _CUDA_VSTD::clamp(nominal_4B_items_per_thread * 4 / int{sizeof(KeyT)}, 1, nominal_4B_items_per_thread);
     using RleSweepPolicyT =
       AgentRlePolicy<96,
                      ITEMS_PER_THREAD,
diff --git a/cub/cub/device/dispatch/tuning/tuning_select_if.cuh b/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
index d3f99ee8ee8..ea2971dd3d2 100644
--- a/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_select_if.cuh
@@ -45,6 +45,8 @@
 #include <cub/util_math.cuh>
 #include <cub/util_type.cuh>
 
+#include <cuda/std/__algorithm_>
+
 CUB_NAMESPACE_BEGIN
 
 namespace detail
@@ -1498,9 +1500,8 @@ struct policy_hub
   struct DefaultPolicy
   {
     static constexpr int nominal_4B_items_per_thread = 10;
-    // TODO(bgruber): use cuda::std::clamp() in C++14
     static constexpr int items_per_thread =
-      CUB_MIN(nominal_4B_items_per_thread, CUB_MAX(1, (nominal_4B_items_per_thread * 4 / sizeof(InputT))));
+      ::cuda::std::clamp(nominal_4B_items_per_thread * 4 / int{sizeof(InputT)}, 1, nominal_4B_items_per_thread);
     using SelectIfPolicyT =
       AgentSelectIfPolicy<128,
                           items_per_thread,
diff --git a/thrust/testing/zip_iterator.cu b/thrust/testing/zip_iterator.cu
index 786b839099b..47b8033bfb1 100644
--- a/thrust/testing/zip_iterator.cu
+++ b/thrust/testing/zip_iterator.cu
@@ -332,25 +332,25 @@ struct TestZipIteratorTransform
     device_vector<T> d_result(n);
 
     // Tuples with 2 elements
-    transform(make_zip_iterator(h_data0.begin(), h_data1.begin()),
-              make_zip_iterator(h_data0.end(), h_data1.end()),
-              h_result.begin(),
-              SumTwoTuple());
-    transform(make_zip_iterator(d_data0.begin(), d_data1.begin()),
-              make_zip_iterator(d_data0.end(), d_data1.end()),
-              d_result.begin(),
-              SumTwoTuple());
+    thrust::transform(make_zip_iterator(h_data0.begin(), h_data1.begin()),
+                      make_zip_iterator(h_data0.end(), h_data1.end()),
+                      h_result.begin(),
+                      SumTwoTuple());
+    thrust::transform(make_zip_iterator(d_data0.begin(), d_data1.begin()),
+                      make_zip_iterator(d_data0.end(), d_data1.end()),
+                      d_result.begin(),
+                      SumTwoTuple());
     ASSERT_EQUAL(h_result, d_result);
 
     // Tuples with 3 elements
-    transform(make_zip_iterator(h_data0.begin(), h_data1.begin(), h_data2.begin()),
-              make_zip_iterator(h_data0.end(), h_data1.end(), h_data2.end()),
-              h_result.begin(),
-              SumThreeTuple());
-    transform(make_zip_iterator(d_data0.begin(), d_data1.begin(), d_data2.begin()),
-              make_zip_iterator(d_data0.end(), d_data1.end(), d_data2.end()),
-              d_result.begin(),
-              SumThreeTuple());
+    thrust::transform(make_zip_iterator(h_data0.begin(), h_data1.begin(), h_data2.begin()),
+                      make_zip_iterator(h_data0.end(), h_data1.end(), h_data2.end()),
+                      h_result.begin(),
+                      SumThreeTuple());
+    thrust::transform(make_zip_iterator(d_data0.begin(), d_data1.begin(), d_data2.begin()),
+                      make_zip_iterator(d_data0.end(), d_data1.end(), d_data2.end()),
+                      d_result.begin(),
+                      SumThreeTuple());
     ASSERT_EQUAL(h_result, d_result);
   }
 };

From cd6a09097f8b92a5f81e625c0ffd2f56fda24d00 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Wed, 26 Feb 2025 16:08:13 +0100
Subject: [PATCH 17/17] Deprecate and replace some Thrust iterator traits
 (#3928)

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..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 = typename thrust::iterator_difference<InputIterator1>::type;
+  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 c86bb22f64e..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 = typename thrust::iterator_difference<Iterator>::type;
+  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 40ad38fcb9a..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 = typename thrust::iterator_difference<Iterator>::type;
+  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 ca69dea2f77..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 = typename thrust::iterator_difference<Iterator>::type;
+  using difference_type = typename cuda::std::iterator_traits<Iterator>::difference_type;
 
   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 47b8033bfb1..00611d23356 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 94dd441c15b..ae049c53e5e 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 12eff7a94a8..8bedfd5d818 100644
--- a/thrust/thrust/system/cuda/detail/core/util.h
+++ b/thrust/thrust/system/cuda/detail/core/util.h
@@ -482,7 +482,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 60a093c05dc..a39118c99c6 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 f0ce7a1c62d..474836f2b57 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 dabafcaaf49..88694e2075d 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 1982ab6c19e..7b0a4d7c22d 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 b1344061495..d7f449ced56 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 216e7996565..25e8a53a76d 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 6edb23cee86..cfe8dc134ee 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