NVIDIA · fbusato · Apr 1, 2026 · Apr 1, 2026 · Apr 1, 2026 · Apr 1, 2026
@@ -16,6 +16,7 @@
 #include <cub/util_temporary_storage.cuh>
 #include <cub/util_type.cuh>
 
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/cstdint>
 #include <cuda/std/memory>
 
@@ -327,7 +328,7 @@ static_assert(device_transform_policy()(detail::current_tuning_cc()) == {9}, "Ho
   build_ptr->cache                      = new transform::cache();
 
   // avoid new and delete which requires the allocated and freed types to match
-  static_assert(std::is_trivially_copyable_v<decltype(policy_sel)>);
+  static_assert(::cuda::is_trivially_copyable_v<decltype(policy_sel)>);
   build_ptr->runtime_policy = std::malloc(sizeof(policy_sel));
   std::memcpy(build_ptr->runtime_policy, &policy_sel, sizeof(policy_sel));
 
@@ -528,7 +529,7 @@ static_assert(device_transform_policy()(detail::current_tuning_cc()) == {12}, "H
   build_ptr->cache                      = new transform::cache();
 
   // avoid new and delete which requires the allocated and freed types to match
-  static_assert(std::is_trivially_copyable_v<decltype(policy_sel)>);
+  static_assert(::cuda::is_trivially_copyable_v<decltype(policy_sel)>);
   build_ptr->runtime_policy = std::malloc(sizeof(policy_sel));
   std::memcpy(build_ptr->runtime_policy, &policy_sel, sizeof(policy_sel));
 

@@ -29,9 +29,9 @@
 #include <cub/util_temporary_storage.cuh>
 #include <cub/warp/warp_reduce.cuh>
 
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__type_traits/conditional.h>
 #include <cuda/std/__type_traits/enable_if.h>
-#include <cuda/std/__type_traits/is_trivially_copyable.h>
 
 #include <nv/target>
 
@@ -473,7 +473,7 @@ using default_no_delay_t             = default_no_delay_constructor_t::delay_t;
 template <class T>
 using default_delay_constructor_t =
   // TODO(bgruber): remove the check for is_primitive<ValueT> in CCCL 4.0
-  ::cuda::std::conditional_t<is_primitive<T>::value || ::cuda::std::is_trivially_copyable_v<T>,
+  ::cuda::std::conditional_t<is_primitive<T>::value || ::cuda::is_trivially_copyable_v<T>,
                              fixed_delay_constructor_t<350, 450>,
                              default_no_delay_constructor_t>;
 
@@ -483,7 +483,7 @@ using default_delay_t = typename default_delay_constructor_t<T>::delay_t;
 template <class KeyT, class ValueT>
 using default_reduce_by_key_delay_constructor_t =
   // TODO(bgruber): remove the check for is_primitive<ValueT> in CCCL 4.0
-  ::cuda::std::conditional_t<(is_primitive<ValueT>::value || ::cuda::std::is_trivially_copyable_v<ValueT>)
+  ::cuda::std::conditional_t<(is_primitive<ValueT>::value || ::cuda::is_trivially_copyable_v<ValueT>)
                                && (sizeof(ValueT) + sizeof(KeyT) < largest_atomic_message_size),
                              reduce_by_key_delay_constructor_t<350, 450>,
                              default_delay_constructor_t<KeyValuePair<KeyT, ValueT>>>;
@@ -582,7 +582,7 @@ _CCCL_HOST_DEVICE _CCCL_FORCEINLINE cudaError_t tile_state_init(
 template <typename T,
           // TODO(bgruber): remove the check for is_primitive<T> in CCCL 4.0
           bool SingleWord = detail::is_primitive<T>::value
-                         || (::cuda::std::is_trivially_copyable_v<T>
+                         || (::cuda::is_trivially_copyable_v<T>
                              && sizeof(T) < detail::largest_atomic_message_size
                              // TODO(bgruber): a power of two size is not strictly necessary, but the implementation
                              // cannot handle it currently. For example, we could support status word + int3.
@@ -957,7 +957,7 @@ struct ScanTileState<T, false>
 template <typename ValueT,
           typename KeyT,
           // TODO(bgruber): remove the check for is_primitive<ValueT> in CCCL 4.0
-          bool SingleWord = (detail::is_primitive<ValueT>::value || ::cuda::std::is_trivially_copyable_v<ValueT>)
+          bool SingleWord = (detail::is_primitive<ValueT>::value || ::cuda::is_trivially_copyable_v<ValueT>)
                          && (sizeof(ValueT) + sizeof(KeyT) < detail::largest_atomic_message_size)>
 struct ReduceByKeyScanTileState;
 

@@ -13,9 +13,9 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__new/device_new.h>
 #include <cuda/std/__type_traits/enable_if.h>
-#include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__utility/forward.h>
 
 CUB_NAMESPACE_BEGIN
@@ -30,14 +30,14 @@ _CCCL_HOST_DEVICE void uninitialized_copy_single(T* ptr, U&& val)
   new (ptr) T(::cuda::std::forward<U>(val));
 }
 #else // ^^^ _CCCL_CUDA_COMPILER(NVHPC) ^^^ / vvv !_CCCL_CUDA_COMPILER(NVHPC) vvv
-template <typename T, typename U, ::cuda::std::enable_if_t<::cuda::std::is_trivially_copyable_v<T>, int> = 0>
+template <typename T, typename U, ::cuda::std::enable_if_t<::cuda::is_trivially_copyable_v<T>, int> = 0>
 _CCCL_HOST_DEVICE void uninitialized_copy_single(T* ptr, U&& val)
 {
   // gevtushenko: placement new should work here as well, but the code generated for copy assignment is sometimes better
   *ptr = ::cuda::std::forward<U>(val);
 }
 
-template <typename T, typename U, ::cuda::std::enable_if_t<!::cuda::std::is_trivially_copyable_v<T>, int> = 0>
+template <typename T, typename U, ::cuda::std::enable_if_t<!::cuda::is_trivially_copyable_v<T>, int> = 0>
 _CCCL_HOST_DEVICE void uninitialized_copy_single(T* ptr, U&& val)
 {
   new (ptr) T(::cuda::std::forward<U>(val));

@@ -23,6 +23,7 @@
 #include <cuda/__functional/operator_properties.h>
 #include <cuda/__memory/is_aligned.h>
 #include <cuda/__ptx/instructions/get_sreg.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__bit/popcount.h>
 #include <cuda/std/__type_traits/underlying_type.h>
 
@@ -75,7 +76,7 @@ storeTileAggregate(tile_state_t<AccumT>* ptrTileStates, scan_state scanState, Ac
   _CCCL_ASSERT(index >= 0 && index < gridDim.x, "Reading out of bounds tile state");
 
   if constexpr (sizeof(tile_state_t<AccumT>) <= cub::detail::warpspeed::max_native_atomic_size()
-                && ::cuda::std::is_trivially_copyable_v<tile_state_t<AccumT>>)
+                && ::cuda::is_trivially_copyable_v<tile_state_t<AccumT>>)
   {
     static_assert(::cuda::is_power_of_two(sizeof(tile_state_t<AccumT>)));
     tile_state_t<AccumT> tmp{scanState, sum};
@@ -103,7 +104,7 @@ _CCCL_DEVICE_API tile_state_t<AccumT> loadTileAggregate(tile_state_t<AccumT>* pt
 
   tile_state_t<AccumT> res;
   if constexpr (sizeof(tile_state_t<AccumT>) <= cub::detail::warpspeed::max_native_atomic_size()
-                && ::cuda::std::is_trivially_copyable_v<tile_state_t<AccumT>>)
+                && ::cuda::is_trivially_copyable_v<tile_state_t<AccumT>>)
   {
     static_assert(::cuda::is_power_of_two(sizeof(tile_state_t<AccumT>)));
 #  if _CCCL_HAS_NV_ATOMIC_BUILTINS()

@@ -18,6 +18,7 @@
 #include <cub/grid/grid_queue.cuh>
 #include <cub/util_arch.cuh>
 
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__numeric/reduce.h>
 
 CUB_NAMESPACE_BEGIN
@@ -80,7 +81,7 @@ struct Transforms
     static_assert(::cuda::std::is_convertible_v<CommonT, int>,
                   "The common type of `LevelT` and `SampleT` must be "
                   "convertible to `int`.");
-    static_assert(::cuda::std::is_trivially_copyable_v<CommonT>,
+    static_assert(::cuda::is_trivially_copyable_v<CommonT>,
                   "The common type of `LevelT` and `SampleT` must be "
                   "trivially copyable.");
 

@@ -24,9 +24,9 @@
 
 #include <cuda/__cmath/ceil_div.h>
 #include <cuda/__device/compute_capability.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__algorithm/clamp.h>
 #include <cuda/std/__host_stdlib/ostream>
-#include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/concepts>
 
 CUB_NAMESPACE_BEGIN
@@ -1657,7 +1657,7 @@ struct policy_selector_from_types
       int{sizeof(AccumT)},
       classify_type<AccumT>,
       is_primitive_v<KeyT>,
-      ::cuda::std::is_trivially_copyable_v<KeyT>,
+      ::cuda::is_trivially_copyable_v<KeyT>,
       is_primitive_v<AccumT>,
       basic_binary_op_t<ReductionOpT>::value}(cc);
   }

@@ -25,6 +25,7 @@
 
 #include <cuda/__cmath/ceil_div.h>
 #include <cuda/__device/compute_capability.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__algorithm/clamp.h>
 #include <cuda/std/__algorithm/max.h>
 #include <cuda/std/__host_stdlib/ostream>
@@ -525,7 +526,7 @@ struct policy_selector_from_types
       int{sizeof(KeyT)},
       classify_type<KeyT>,
       is_primitive_v<LengthT>,
-      ::cuda::std::is_trivially_copyable_v<LengthT>,
+      ::cuda::is_trivially_copyable_v<LengthT>,
       is_primitive_v<KeyT>};
     return selector(cc);
   }

@@ -24,6 +24,7 @@
 #include <cub/util_device.cuh>
 
 #include <cuda/__device/compute_capability.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__algorithm/clamp.h>
 #include <cuda/std/__host_stdlib/ostream>
 #include <cuda/std/concepts>
@@ -612,7 +613,7 @@ struct policy_selector_from_types
       classify_type<KeyT>,
       is_primitive_v<LengthT>,
       is_primitive_v<KeyT>,
-      ::cuda::std::is_trivially_copyable_v<KeyT>};
+      ::cuda::is_trivially_copyable_v<KeyT>};
     return selector(cc);
   }
 };

@@ -31,6 +31,7 @@
 #include <thrust/type_traits/is_contiguous_iterator.h>
 
 #include <cuda/__device/compute_capability.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__algorithm/max.h>
 #include <cuda/std/__functional/invoke.h>
 #include <cuda/std/__functional/operations.h>
@@ -1451,8 +1452,8 @@ struct policy_selector_from_types
       classify_op<ScanOpT>,
       THRUST_NS_QUALIFIER::is_contiguous_iterator_v<InputIteratorT>,
       THRUST_NS_QUALIFIER::is_contiguous_iterator_v<OutputIteratorT>,
-      ::cuda::std::is_trivially_copyable_v<InputValueT>,
-      ::cuda::std::is_trivially_copyable_v<OutputValueT>,
+      ::cuda::is_trivially_copyable_v<InputValueT>,
+      ::cuda::is_trivially_copyable_v<OutputValueT>,
       ::cuda::std::is_default_constructible_v<OutputValueT>,
       accum_is_primitive_or_trivially_copy_constructible,
       benchmark_match};

@@ -26,9 +26,9 @@
 #include <cub/util_type.cuh>
 
 #include <cuda/__device/compute_capability.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__algorithm/max.h>
 #include <cuda/std/__host_stdlib/ostream>
-#include <cuda/std/__type_traits/is_trivially_copyable.h>
 
 CUB_NAMESPACE_BEGIN
 
@@ -1951,7 +1951,7 @@ struct policy_selector_from_types
       static_cast<int>(sizeof(ValueT)),
       static_cast<int>(sizeof(AccumT)),
       is_primitive<ValueT>::value,
-      ::cuda::std::is_trivially_copyable_v<ValueT>,
+      ::cuda::is_trivially_copyable_v<ValueT>,
       classify_type<KeyT>,
       classify_type<ValueT>,
       classify_type<AccumT>,

@@ -27,8 +27,8 @@
 #  include <cuda/__runtime/api_wrapper.h>
 #  include <cuda/__stream/invalid_stream.h>
 #  include <cuda/__stream/stream_ref.h>
+#  include <cuda/__type_traits/is_trivially_copyable.h>
 #  include <cuda/__utility/no_init.h>
-#  include <cuda/std/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/__utility/exchange.h>
 #  include <cuda/std/__utility/move.h>
 #  include <cuda/std/cstddef>
@@ -68,7 +68,7 @@ namespace cuda::experimental
 template <class _Tp>
 class graph_buffer
 {
-  static_assert(::cuda::std::is_trivially_copyable_v<_Tp>, "graph_buffer requires T to be trivially copyable.");
+  static_assert(::cuda::is_trivially_copyable_v<_Tp>, "graph_buffer requires T to be trivially copyable.");
 
 public:
   using value_type      = _Tp;

@@ -33,6 +33,7 @@
 #  include <cuda/__mdspan/host_device_mdspan.h>
 #  include <cuda/__mdspan/traits.h>
 #  include <cuda/__stream/stream_ref.h>
+#  include <cuda/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/__algorithm/max.h>
 #  include <cuda/std/__functional/identity.h>
 #  include <cuda/std/__host_stdlib/stdexcept>
@@ -43,7 +44,6 @@
 #  include <cuda/std/__type_traits/is_const.h>
 #  include <cuda/std/__type_traits/is_convertible.h>
 #  include <cuda/std/__type_traits/is_same.h>
-#  include <cuda/std/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/__type_traits/remove_cvref.h>
 
 #  include <cuda/experimental/__copy/copy_contiguous.cuh>
@@ -125,7 +125,7 @@ _CCCL_HOST_API void copy(::cuda::device_mdspan<_TpIn, _ExtentsIn, _LayoutPolicyI
     && ::cuda::std::is_convertible_v<_AccessorPolicyOut, __default_accessor_out>;
   constexpr bool __are_byte_copyable =
     ::cuda::std::is_same_v<::cuda::std::remove_cv_t<_TpIn>, ::cuda::std::remove_cv_t<_TpOut>>
-    && ::cuda::std::is_trivially_copyable_v<_TpIn> //
+    && ::cuda::is_trivially_copyable_v<_TpIn> //
     && __have_default_accessors;
 
   if (__tensor_size == 1 && __are_byte_copyable)

@@ -27,6 +27,7 @@
 #  include <cuda/__mdspan/host_device_mdspan.h>
 #  include <cuda/__mdspan/traits.h>
 #  include <cuda/__stream/stream_ref.h>
+#  include <cuda/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/__algorithm/max.h>
 #  include <cuda/std/__cstddef/types.h>
 #  include <cuda/std/__host_stdlib/stdexcept>
@@ -37,7 +38,6 @@
 #  include <cuda/std/__type_traits/is_const.h>
 #  include <cuda/std/__type_traits/is_convertible.h>
 #  include <cuda/std/__type_traits/is_same.h>
-#  include <cuda/std/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/__type_traits/remove_cv.h>
 
 #  include <cuda/experimental/__copy_bytes/memcpy_batch_tiles.cuh>
@@ -74,7 +74,7 @@ _CCCL_HOST_API void __copy_bytes_impl(
   namespace cudax = ::cuda::experimental;
   static_assert(::cuda::std::is_same_v<::cuda::std::remove_cv_t<_TpIn>, ::cuda::std::remove_cv_t<_TpOut>>,
                 "cudax::copy_bytes: TpIn and TpOut must be the same type");
-  static_assert(::cuda::std::is_trivially_copyable_v<_TpIn>, "TpIn must be trivially copyable");
+  static_assert(::cuda::is_trivially_copyable_v<_TpIn>, "TpIn must be trivially copyable");
   static_assert(!::cuda::std::is_const_v<_TpOut>, "TpOut must not be const");
   static_assert(::cuda::__is_cuda_mdspan_layout_v<_LayoutPolicyIn>,
                 "cudax::copy_bytes: LayoutPolicyIn must be a predefined layout policy");

@@ -26,11 +26,11 @@
 #  include <cuda/__algorithm/common.h>
 #  include <cuda/__stream/launch_transform.h>
 #  include <cuda/__stream/stream_ref.h>
+#  include <cuda/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/__concepts/concept_macros.h>
 #  include <cuda/std/__exception/exception_macros.h>
 #  include <cuda/std/__host_stdlib/stdexcept>
 #  include <cuda/std/__type_traits/is_const.h>
-#  include <cuda/std/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/cstddef>
 #  include <cuda/std/span>
 
@@ -47,7 +47,7 @@ _CCCL_HOST_API graph_node_ref
 __copy_bytes_graph_impl(path_builder& __pb, ::cuda::std::span<_SrcTy> __src, ::cuda::std::span<_DstTy> __dst)
 {
   static_assert(!::cuda::std::is_const_v<_DstTy>, "Copy destination can't be const");
-  static_assert(::cuda::std::is_trivially_copyable_v<_SrcTy> && ::cuda::std::is_trivially_copyable_v<_DstTy>,
+  static_assert(::cuda::is_trivially_copyable_v<_SrcTy> && ::cuda::is_trivially_copyable_v<_DstTy>,
                 "Copy source and destination element types must be trivially copyable");
 
   if (__src.size_bytes() > __dst.size_bytes())

@@ -26,11 +26,11 @@
 #  include <cuda/__algorithm/common.h>
 #  include <cuda/__stream/launch_transform.h>
 #  include <cuda/__stream/stream_ref.h>
+#  include <cuda/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/__concepts/concept_macros.h>
 #  include <cuda/std/__exception/exception_macros.h>
 #  include <cuda/std/__host_stdlib/stdexcept>
 #  include <cuda/std/__type_traits/is_const.h>
-#  include <cuda/std/__type_traits/is_trivially_copyable.h>
 #  include <cuda/std/cstdint>
 #  include <cuda/std/span>
 
@@ -48,8 +48,7 @@ _CCCL_HOST_API graph_node_ref
 __fill_bytes_graph_impl(path_builder& __pb, ::cuda::std::span<_DstTy, _DstSize> __dst, ::cuda::std::uint8_t __value)
 {
   static_assert(!::cuda::std::is_const_v<_DstTy>, "Fill destination can't be const");
-  static_assert(::cuda::std::is_trivially_copyable_v<_DstTy>,
-                "Fill destination element type must be trivially copyable");
+  static_assert(::cuda::is_trivially_copyable_v<_DstTy>, "Fill destination element type must be trivially copyable");
 
   auto __deps = __pb.get_dependencies();
   ::CUgraphNodeParams __params{};

@@ -25,8 +25,8 @@
 #include <cuda/__driver/driver_api.h>
 #include <cuda/__memory/address_space.h>
 #include <cuda/__runtime/api_wrapper.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__type_traits/always_false.h>
-#include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/string_view>
 
@@ -51,7 +51,7 @@ class kernel_ref
 template <class... _Args>
 class kernel_ref<void(_Args...)>
 {
-  static_assert((true && ... && ::cuda::std::is_trivially_copyable_v<_Args>),
+  static_assert((true && ... && ::cuda::is_trivially_copyable_v<_Args>),
                 "All kernel_ref argument types must be trivially copyable.");
 
 public:

@@ -8,6 +8,7 @@ Type traits
    :maxdepth: 1
 
    type_traits/is_floating_point
+   type_traits/is_trivially_copyable
    type_traits/is_bitwise_comparable
    type_traits/vector_types
 
@@ -30,6 +31,11 @@ Type traits
      - CCCL 3.3.0
      - CUDA 13.3
 
+   * - :ref:`cuda::is_trivially_copyable <libcudacxx-extended-api-type_traits-is_trivially_copyable>`
+     - Relaxed trivially copyable check including extended floating-point types
+     - CCCL 3.4.0
+     - CUDA 13.4
+
    * - :ref:`cuda::is_bitwise_comparable <libcudacxx-extended-api-type_traits-is_bitwise_comparable>`
      - User-specializable bitwise comparability check
      - CCCL 3.4.0