diff --git a/docs/libcudacxx/extended_api/type_traits.rst b/docs/libcudacxx/extended_api/type_traits.rst
index 8588b77c375..36148f5af7a 100644
--- a/docs/libcudacxx/extended_api/type_traits.rst
+++ b/docs/libcudacxx/extended_api/type_traits.rst
@@ -8,6 +8,7 @@ Type traits
    :maxdepth: 1
 
    type_traits/is_floating_point
+   type_traits/is_trivially_copyable
    type_traits/vector_types
 
 .. list-table::
@@ -28,3 +29,8 @@ Type traits
      - Type traits for CUDA vector types
      - 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
diff --git a/docs/libcudacxx/extended_api/type_traits/is_trivially_copyable.rst b/docs/libcudacxx/extended_api/type_traits/is_trivially_copyable.rst
new file mode 100644
index 00000000000..f6164f04948
--- /dev/null
+++ b/docs/libcudacxx/extended_api/type_traits/is_trivially_copyable.rst
@@ -0,0 +1,68 @@
+.. _libcudacxx-extended-api-type_traits-is_trivially_copyable:
+
+``cuda::is_trivially_copyable``
+=======================================
+
+Defined in the ``<cuda/type_traits>`` header.
+
+.. code:: cuda
+
+   namespace cuda {
+
+   template <typename T>
+   constexpr bool is_trivially_copyable_v = /* see below */;
+
+   template <typename T>
+   using is_trivially_copyable = cuda::std::bool_constant<is_trivially_copyable_v<T>>;
+
+   } // namespace cuda
+
+``cuda::is_trivially_copyable_v<T>`` is a variable template that extends ``cuda::std::is_trivially_copyable`` to also recognize CUDA extended floating-point scalar and vector types as trivially copyable.
+
+A type ``T`` satisfies ``cuda::is_trivially_copyable`` if any of the following holds:
+
+- ``T`` is trivially copyable.
+- ``T`` is an extended floating-point scalar type (e.g. ``__half``, ``__nv_bfloat16``, ``__nv_fp8_e4m3``).
+- ``T`` is an extended floating-point vector type (e.g. ``__half2``, ``__nv_bfloat162``, ``__nv_fp8x2_e4m3``).
+
+The trait also propagates through composite types:
+
+- C-style arrays: ``T[N]`` and ``T[]`` are trivially copyable when ``T`` is.
+- ``cuda::std::array<T, N>``: trivially copyable when ``T`` is also trivially copyable.
+- ``cuda::std::pair<T1, T2>``: trivially copyable when both ``T1`` and ``T2`` are.
+- ``cuda::std::tuple<Ts...>``: trivially copyable when all ``Ts...`` are.
+- ``cuda::std::complex<T>``: trivially copyable when ``T`` is.
+- ``cuda::complex<T>``: trivially copyable when ``T`` is.
+- Aggregates: trivially copyable when all members are.
+
+``const`` qualification is handled transparently, while ``volatile`` is compiler dependent.
+
+Examples
+--------
+
+.. code:: cuda
+
+   #include <cuda/type_traits>
+   #include <cuda/std/array>
+   #include <cuda/std/tuple>
+   #include <cuda/std/utility>
+
+   #include <cuda_fp16.h>
+
+   // Standard trivially copyable types
+   static_assert(cuda::is_trivially_copyable_v<int>);
+   static_assert(cuda::is_trivially_copyable_v<float>);
+
+   // Extended floating-point types
+   static_assert(cuda::is_trivially_copyable_v<__half>);
+   static_assert(cuda::is_trivially_copyable_v<__nv_bfloat16>);
+   static_assert(cuda::is_trivially_copyable_v<__half2>);
+   static_assert(cuda::is_trivially_copyable_v<cuda::std::complex<__half>>);
+   static_assert(cuda::is_trivially_copyable_v<cuda::complex<__half>>);
+
+   // Composite types containing extended floating-point types
+   static_assert(cuda::is_trivially_copyable_v<__half[4]>);
+   static_assert(cuda::is_trivially_copyable_v<cuda::std::array<__half, 4>>);
+   static_assert(cuda::is_trivially_copyable_v<cuda::std::pair<__half, __half>>);
+   static_assert(cuda::is_trivially_copyable_v<cuda::std::tuple<__half, __half>>);
+   static_assert(cuda::is_trivially_copyable_v<cuda::std::pair<__half, int>>);
diff --git a/docs/libcudacxx/standard_api/numerics_library/bit.rst b/docs/libcudacxx/standard_api/numerics_library/bit.rst
index 9904db7352e..44079a4d81a 100644
--- a/docs/libcudacxx/standard_api/numerics_library/bit.rst
+++ b/docs/libcudacxx/standard_api/numerics_library/bit.rst
@@ -3,6 +3,16 @@
 ``<cuda/std/bit>``
 ==================
 
+``cuda::std::bit_cast``
+-----------------------
+
+``cuda::std::bit_cast`` extended the standard ``std::bit_cast`` to also recognize CUDA extended floating-point scalar and vector types as trivially copyable.
+
+**Limitations**
+
+- The function can be used in ``constexpr`` contexts only when the source and destination types are trivially copyable.
+- The function cannot be used in ``constexpr`` contexts with MSVC <= 19.25 and GCC <= 10.
+
 CUDA Performance Considerations
 -------------------------------
 
diff --git a/libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h b/libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h
new file mode 100644
index 00000000000..61f6b9bc906
--- /dev/null
+++ b/libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h
@@ -0,0 +1,99 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CUDA__TYPE_TRAITS_IS_TRIVIALLY_COPYABLE_H
+#define __CUDA__TYPE_TRAITS_IS_TRIVIALLY_COPYABLE_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/__fwd/complex.h>
+#include <cuda/__type_traits/is_vector_type.h>
+#include <cuda/std/__cstddef/types.h>
+#include <cuda/std/__fwd/array.h>
+#include <cuda/std/__fwd/complex.h>
+#include <cuda/std/__fwd/pair.h>
+#include <cuda/std/__fwd/tuple.h>
+#include <cuda/std/__type_traits/aggregate_members.h>
+#include <cuda/std/__type_traits/enable_if.h>
+#include <cuda/std/__type_traits/integral_constant.h>
+#include <cuda/std/__type_traits/is_extended_floating_point.h>
+#include <cuda/std/__type_traits/is_trivially_copyable.h>
+#include <cuda/std/__type_traits/remove_const.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA
+
+template <typename _Tp, typename = void>
+inline constexpr bool __is_aggregate_trivially_copyable_v = false;
+
+//! Users are allowed to specialize this variable template for their own types
+template <typename _Tp>
+inline constexpr bool __is_trivially_copyable_v =
+  ::cuda::std::is_trivially_copyable_v<_Tp> || ::cuda::std::__is_extended_floating_point_v<_Tp>
+#if _CCCL_HAS_CTK()
+  || ::cuda::is_extended_fp_vector_type_v<_Tp>
+#endif // _CCCL_HAS_CTK()
+  || __is_aggregate_trivially_copyable_v<_Tp>;
+
+template <typename _Tp>
+inline constexpr bool __is_trivially_copyable_v<_Tp[]> = __is_trivially_copyable_v<_Tp>;
+
+template <typename _Tp, ::cuda::std::size_t _Size>
+inline constexpr bool __is_trivially_copyable_v<_Tp[_Size]> = __is_trivially_copyable_v<_Tp>;
+
+template <typename _Tp, ::cuda::std::size_t _Size>
+inline constexpr bool __is_trivially_copyable_v<::cuda::std::array<_Tp, _Size>> = __is_trivially_copyable_v<_Tp>;
+
+template <typename _T1, typename _T2>
+inline constexpr bool __is_trivially_copyable_v<::cuda::std::pair<_T1, _T2>> =
+  __is_trivially_copyable_v<_T1> && __is_trivially_copyable_v<_T2>;
+
+template <typename... _Ts>
+inline constexpr bool __is_trivially_copyable_v<::cuda::std::tuple<_Ts...>> = (__is_trivially_copyable_v<_Ts> && ...);
+
+template <typename _Tp>
+inline constexpr bool __is_trivially_copyable_v<::cuda::std::complex<_Tp>> = __is_trivially_copyable_v<_Tp>;
+
+template <typename _Tp>
+inline constexpr bool __is_trivially_copyable_v<::cuda::complex<_Tp>> = __is_trivially_copyable_v<_Tp>;
+
+// if all the previous conditions fail, check if the type is an aggregate and all its members are trivially copyable
+template <typename _Tp>
+using __is_trivially_copyable_callable = ::cuda::std::bool_constant<__is_trivially_copyable_v<_Tp>>;
+
+template <typename _Tp>
+inline constexpr bool
+  __is_aggregate_trivially_copyable_v<_Tp, ::cuda::std::enable_if_t<::cuda::std::is_aggregate_v<_Tp>>> =
+    ::cuda::std::__aggregate_all_of<__is_trivially_copyable_callable, _Tp>::value;
+
+//----------------------------------------------------------------------------------------------------------------------
+// public traits
+
+template <typename _Tp>
+inline constexpr bool is_trivially_copyable_v = __is_trivially_copyable_v<::cuda::std::remove_const_t<_Tp>>;
+
+// defined as alias so users cannot specialize it (they should specialize the variable template instead)
+template <typename _Tp>
+using is_trivially_copyable = ::cuda::std::bool_constant<is_trivially_copyable_v<_Tp>>;
+
+_CCCL_END_NAMESPACE_CUDA
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // __CUDA__TYPE_TRAITS_IS_TRIVIALLY_COPYABLE_H
diff --git a/libcudacxx/include/cuda/std/__bit/bit_cast.h b/libcudacxx/include/cuda/std/__bit/bit_cast.h
index e7886ad4976..b17e02bdb40 100644
--- a/libcudacxx/include/cuda/std/__bit/bit_cast.h
+++ b/libcudacxx/include/cuda/std/__bit/bit_cast.h
@@ -4,7 +4,7 @@
 // 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.
+// SPDX-FileCopyrightText: Copyright (c) 2024-26 NVIDIA CORPORATION & AFFILIATES.
 //
 //===----------------------------------------------------------------------===//
 
@@ -21,10 +21,9 @@
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/std/__type_traits/enable_if.h>
-#include <cuda/std/__type_traits/is_extended_floating_point.h>
-#include <cuda/std/__type_traits/is_trivially_copyable.h>
-#include <cuda/std/__type_traits/is_trivially_default_constructible.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
+#include <cuda/std/__concepts/concept_macros.h>
+#include <cuda/std/__concepts/constructible.h>
 #include <cuda/std/cstring>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -42,37 +41,43 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD
 #else // ^^^ _CCCL_BUILTIN_BIT_CAST ^^^ / vvv !_CCCL_BUILTIN_BIT_CAST vvv
 #  define _CCCL_CONSTEXPR_BIT_CAST
 #  define _CCCL_HAS_CONSTEXPR_BIT_CAST() 0
-#  if _CCCL_COMPILER(GCC, >=, 8)
-// GCC starting with GCC8 warns about our extended floating point types having protected data members
+#endif // !_CCCL_BUILTIN_BIT_CAST
+
+#if _CCCL_COMPILER(GCC, >=, 8)
 _CCCL_DIAG_PUSH
 _CCCL_DIAG_SUPPRESS_GCC("-Wclass-memaccess")
-#  endif // _CCCL_COMPILER(GCC, >=, 8)
-#endif // !_CCCL_BUILTIN_BIT_CAST
+#endif // _CCCL_COMPILER(GCC, >=, 8)
 
-template <class _To,
-          class _From,
-          enable_if_t<(sizeof(_To) == sizeof(_From)), int>                                          = 0,
-          enable_if_t<is_trivially_copyable_v<_To> || __is_extended_floating_point_v<_To>, int>     = 0,
-          enable_if_t<is_trivially_copyable_v<_From> || __is_extended_floating_point_v<_From>, int> = 0>
-[[nodiscard]] _CCCL_API inline _CCCL_CONSTEXPR_BIT_CAST _To bit_cast(const _From& __from) noexcept
+template <class _To, class _From>
+[[nodiscard]] _CCCL_API inline _To __bit_cast_memcpy(const _From& __from) noexcept
 {
-#if defined(_CCCL_BUILTIN_BIT_CAST)
-  return _CCCL_BUILTIN_BIT_CAST(_To, __from);
-#else // ^^^ _CCCL_BUILTIN_BIT_CAST ^^^ / vvv !_CCCL_BUILTIN_BIT_CAST vvv
-  static_assert(is_trivially_default_constructible_v<_To>,
-                "The compiler does not support __builtin_bit_cast, so bit_cast additionally requires the destination "
-                "type to be trivially constructible");
+  static_assert(::cuda::std::default_initializable<_To>,
+                "bit_cast memcpy fallback requires the destination type to be default initializable");
   _To __temp;
   ::cuda::std::memcpy(&__temp, &__from, sizeof(_To));
   return __temp;
-#endif // !_CCCL_BUILTIN_BIT_CAST
 }
 
-#if !defined(_CCCL_BUILTIN_BIT_CAST)
-#  if _CCCL_COMPILER(GCC, >=, 8)
+#if _CCCL_COMPILER(GCC, >=, 8)
 _CCCL_DIAG_POP
-#  endif // _CCCL_COMPILER(GCC, >=, 8)
-#endif // !_CCCL_BUILTIN_BIT_CAST
+#endif // _CCCL_COMPILER(GCC, >=, 8)
+
+_CCCL_TEMPLATE(class _To, class _From)
+_CCCL_REQUIRES((sizeof(_To) == sizeof(_From)) _CCCL_AND(::cuda::is_trivially_copyable_v<_To>)
+                 _CCCL_AND(::cuda::is_trivially_copyable_v<_From>))
+[[nodiscard]] _CCCL_API inline _CCCL_CONSTEXPR_BIT_CAST _To bit_cast(const _From& __from) noexcept
+{
+#if defined(_CCCL_BUILTIN_BIT_CAST)
+  if constexpr (::cuda::std::is_trivially_copyable_v<_To> && ::cuda::std::is_trivially_copyable_v<_From>)
+  {
+    return _CCCL_BUILTIN_BIT_CAST(_To, __from);
+  }
+  else
+#endif // _CCCL_BUILTIN_BIT_CAST
+  {
+    return ::cuda::std::__bit_cast_memcpy<_To>(__from);
+  }
+}
 
 _CCCL_END_NAMESPACE_CUDA_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/aggregate_members.h b/libcudacxx/include/cuda/std/__type_traits/aggregate_members.h
new file mode 100644
index 00000000000..7ccdab175ea
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__type_traits/aggregate_members.h
@@ -0,0 +1,163 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___TYPE_TRAITS_AGGREGATE_MEMBERS_H
+#define _CUDA_STD___TYPE_TRAITS_AGGREGATE_MEMBERS_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/__type_traits/enable_if.h>
+#include <cuda/std/__type_traits/integral_constant.h>
+#include <cuda/std/__type_traits/is_aggregate.h>
+#include <cuda/std/__type_traits/is_empty.h>
+#include <cuda/std/__type_traits/remove_cvref.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_DIAG_PUSH
+_CCCL_DIAG_SUPPRESS_CLANG("-Wmissing-field-initializers")
+
+_CCCL_BEGIN_NAMESPACE_CUDA_STD
+
+#if defined(_CCCL_BUILTIN_STRUCTURED_BINDING_SIZE)
+
+// Returns the number of aggregate members, or `-1` if the type is not an aggregate.
+template <typename _Tp, ::cuda::std::enable_if_t<_CCCL_BUILTIN_STRUCTURED_BINDING_SIZE(_Tp) >= 0, int> = 0>
+inline constexpr int __aggregate_arity_v = _CCCL_BUILTIN_STRUCTURED_BINDING_SIZE(_Tp);
+
+#else // ^^^ _CCCL_BUILTIN_STRUCTURED_BINDING_SIZE ^^^ / !_CCCL_BUILTIN_STRUCTURED_BINDING_SIZE vvv
+
+// provide a generic way to initialize an aggregate member
+struct __any_aggregate_member
+{
+  template <typename _Tp>
+  _CCCL_API constexpr operator _Tp&&() const;
+};
+
+template <typename _Tp, bool = is_aggregate_v<_Tp>>
+struct __aggregate_arity_impl
+{
+  template <typename... _Args,
+            typename _Up   = _Tp,
+            typename       = decltype(_Up{(_Args{})...}), // SFINAE on args number
+            typename _Self = __aggregate_arity_impl>
+  _CCCL_API auto operator()(_Args... __args) -> decltype(_Self{}(__args..., __any_aggregate_member{}));
+
+  template <typename... _Args>
+  _CCCL_API auto operator()(_Args...) const -> char (*)[sizeof...(_Args) + 1]; // return the number of members + 1
+};
+
+// T is not an aggregate, return 1
+template <typename _Tp>
+struct __aggregate_arity_impl<_Tp, false>
+{
+  _CCCL_API auto operator()() const -> char*;
+};
+
+// Returns the number of aggregate members, or `-1` if the type is not an aggregate.
+template <typename _Tp>
+inline constexpr int __aggregate_arity_v = int{sizeof(*__aggregate_arity_impl<_Tp>{}())} - 2;
+
+#endif // ^^^ !_CCCL_BUILTIN_STRUCTURED_BINDING_SIZE ^^^
+
+// Apply a Predicate to every aggregate member
+
+// provide a generic way to initialize an aggregate member but only if the Predicate is true
+template <template <typename> class _Predicate>
+struct __aggregate_member_if
+{
+  template <typename _Tp, typename = enable_if_t<_Predicate<remove_cvref_t<_Tp>>::value>>
+  _CCCL_API constexpr operator _Tp&&() const;
+};
+
+template <int _Arity>
+struct __aggregate_all_of_fn;
+
+// T has 0 members, return true
+template <>
+struct __aggregate_all_of_fn<0>
+{
+  template <template <typename> class _Predicate, typename _Tp>
+  _CCCL_API static auto __call(int) -> true_type;
+};
+
+#define _CCCL_AGGR_PROBE(_POS) \
+  , __aggregate_member_if<_Predicate> {}
+
+// T has N members, return true if the Predicate is true for all members (recursively)
+#define _CCCL_AGGREGATE_ALL_OF_CASE(_NP)                                                                               \
+  template <>                                                                                                          \
+  struct __aggregate_all_of_fn<1 + _NP>                                                                                \
+  {                                                                                                                    \
+    template <template <typename> class _Predicate,                                                                    \
+              typename _Tp,                                                                                            \
+              typename _Up = _Tp,                                                                                      \
+              typename = decltype(_Up{__aggregate_member_if<_Predicate>{} _CCCL_PP_REPEAT(_NP, _CCCL_AGGR_PROBE, 0)})> \
+    _CCCL_API static auto __call(int) -> true_type;                                                                    \
+                                                                                                                       \
+    template <template <typename> class _Predicate, typename _Tp>                                                      \
+    _CCCL_API static auto __call(...) -> false_type;                                                                   \
+  }
+
+inline constexpr int __aggregate_max_arity = 16;
+
+_CCCL_AGGREGATE_ALL_OF_CASE(0);
+_CCCL_AGGREGATE_ALL_OF_CASE(1);
+_CCCL_AGGREGATE_ALL_OF_CASE(2);
+_CCCL_AGGREGATE_ALL_OF_CASE(3);
+_CCCL_AGGREGATE_ALL_OF_CASE(4);
+_CCCL_AGGREGATE_ALL_OF_CASE(5);
+_CCCL_AGGREGATE_ALL_OF_CASE(6);
+_CCCL_AGGREGATE_ALL_OF_CASE(7);
+_CCCL_AGGREGATE_ALL_OF_CASE(8);
+_CCCL_AGGREGATE_ALL_OF_CASE(9);
+_CCCL_AGGREGATE_ALL_OF_CASE(10);
+_CCCL_AGGREGATE_ALL_OF_CASE(11);
+_CCCL_AGGREGATE_ALL_OF_CASE(12);
+_CCCL_AGGREGATE_ALL_OF_CASE(13);
+_CCCL_AGGREGATE_ALL_OF_CASE(14);
+_CCCL_AGGREGATE_ALL_OF_CASE(15);
+
+#undef _CCCL_AGGREGATE_ALL_OF_CASE
+#undef _CCCL_AGGR_PROBE
+
+// return true if
+// - T is an aggregate
+// - T has a number of members between 0 and __aggregate_max_arity
+// - T is not empty
+template <template <typename> class _Predicate,
+          typename _Tp,
+          bool = is_aggregate_v<_Tp> && (__aggregate_arity_v<_Tp> >= 0)
+              && (__aggregate_arity_v<_Tp> <= __aggregate_max_arity)
+              && ((__aggregate_arity_v<_Tp> != 0) || is_empty_v<_Tp>)>
+struct __aggregate_all_of : false_type
+{};
+
+// Applies a Predicate to every member reachable by aggregate initialization
+template <template <typename> class _Predicate, typename _Tp>
+struct __aggregate_all_of<_Predicate, _Tp, true>
+    : decltype(__aggregate_all_of_fn<__aggregate_arity_v<_Tp>>::template __call<_Predicate, _Tp>(0))
+{};
+
+_CCCL_END_NAMESPACE_CUDA_STD
+
+_CCCL_DIAG_POP
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // _CUDA_STD___TYPE_TRAITS_AGGREGATE_MEMBERS_H
diff --git a/libcudacxx/include/cuda/std/__type_traits/aggregate_members_all_of.h b/libcudacxx/include/cuda/std/__type_traits/aggregate_members_all_of.h
new file mode 100644
index 00000000000..321ced87562
--- /dev/null
+++ b/libcudacxx/include/cuda/std/__type_traits/aggregate_members_all_of.h
@@ -0,0 +1,160 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_STD___TYPE_TRAITS_AGGREGATE_MEMBERS_ALL_OF_H
+#define _CUDA_STD___TYPE_TRAITS_AGGREGATE_MEMBERS_ALL_OF_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/__type_traits/enable_if.h>
+#include <cuda/std/__type_traits/integral_constant.h>
+#include <cuda/std/__type_traits/is_aggregate.h>
+#include <cuda/std/__type_traits/is_empty.h>
+#include <cuda/std/__type_traits/remove_cvref.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_DIAG_PUSH
+_CCCL_DIAG_SUPPRESS_CLANG("-Wmissing-field-initializers")
+_CCCL_DIAG_SUPPRESS_CLANG("-Wmissing-braces")
+
+_CCCL_BEGIN_NAMESPACE_CUDA_STD
+
+// NOTE: __builtin_structured_binding_size cannot be used here. It returns the structural member count (e.g., 1 for a
+// struct with an int[4] member), while the SFINAE-based arity detection below counts the brace-elided elements (4 in
+// that case). Since __aggregate_all_of also relies on brace-elided aggregate initialization to apply the predicate, the
+// arity must match the brace-elided count. Using the builtin's structural count would cause __aggregate_all_of to
+// silently skip unchecked members in aggregates with arrays or nested aggregate members.
+
+// provide a generic way to initialize an aggregate member
+struct __any_aggregate_member
+{
+  template <typename _Tp>
+  _CCCL_API constexpr operator _Tp&&() const;
+};
+
+template <typename _Tp, bool = is_aggregate_v<_Tp>>
+struct __aggregate_arity_impl
+{
+  template <typename... _Args,
+            typename _Up   = _Tp,
+            typename       = decltype(_Up{(_Args{})...}), // SFINAE on args number
+            typename _Self = __aggregate_arity_impl>
+  _CCCL_API auto operator()(_Args... __args) -> decltype(_Self{}(__args..., __any_aggregate_member{}));
+
+  template <typename... _Args>
+  _CCCL_API auto operator()(_Args...) const -> char (*)[sizeof...(_Args) + 1]; // return the number of members + 1
+};
+
+// T is not an aggregate, return 1
+template <typename _Tp>
+struct __aggregate_arity_impl<_Tp, false>
+{
+  _CCCL_API auto operator()() const -> char*;
+};
+
+// Returns the number of aggregate members, or `-1` if the type is not an aggregate.
+template <typename _Tp>
+inline constexpr int __aggregate_arity_v = int{sizeof(*__aggregate_arity_impl<_Tp>{}())} - 2;
+
+// Apply a Predicate to every aggregate member
+
+// provide a generic way to initialize an aggregate member but only if the Predicate is true
+template <template <typename> class _Predicate>
+struct __aggregate_member_if
+{
+  template <typename _Tp, typename = enable_if_t<_Predicate<remove_cvref_t<_Tp>>::value>>
+  _CCCL_API constexpr operator _Tp&&() const;
+};
+
+template <int _Arity>
+struct __aggregate_all_of_fn;
+
+// T has 0 members, return true
+template <>
+struct __aggregate_all_of_fn<0>
+{
+  template <template <typename> class _Predicate, typename _Tp>
+  _CCCL_API static auto __call(int) -> true_type;
+};
+
+#define _CCCL_AGGR_PROBE(_POS) \
+  , __aggregate_member_if<_Predicate> {}
+
+// T has N members, return true if the Predicate is true for all members (recursively)
+#define _CCCL_AGGREGATE_ALL_OF_CASE(_NP)                                                                               \
+  template <>                                                                                                          \
+  struct __aggregate_all_of_fn<1 + _NP>                                                                                \
+  {                                                                                                                    \
+    template <template <typename> class _Predicate,                                                                    \
+              typename _Tp,                                                                                            \
+              typename _Up = _Tp,                                                                                      \
+              typename = decltype(_Up{__aggregate_member_if<_Predicate>{} _CCCL_PP_REPEAT(_NP, _CCCL_AGGR_PROBE, 0)})> \
+    _CCCL_API static auto __call(int) -> true_type;                                                                    \
+                                                                                                                       \
+    template <template <typename> class _Predicate, typename _Tp>                                                      \
+    _CCCL_API static auto __call(...) -> false_type;                                                                   \
+  }
+
+inline constexpr int __aggregate_max_arity = 16;
+
+_CCCL_AGGREGATE_ALL_OF_CASE(0);
+_CCCL_AGGREGATE_ALL_OF_CASE(1);
+_CCCL_AGGREGATE_ALL_OF_CASE(2);
+_CCCL_AGGREGATE_ALL_OF_CASE(3);
+_CCCL_AGGREGATE_ALL_OF_CASE(4);
+_CCCL_AGGREGATE_ALL_OF_CASE(5);
+_CCCL_AGGREGATE_ALL_OF_CASE(6);
+_CCCL_AGGREGATE_ALL_OF_CASE(7);
+_CCCL_AGGREGATE_ALL_OF_CASE(8);
+_CCCL_AGGREGATE_ALL_OF_CASE(9);
+_CCCL_AGGREGATE_ALL_OF_CASE(10);
+_CCCL_AGGREGATE_ALL_OF_CASE(11);
+_CCCL_AGGREGATE_ALL_OF_CASE(12);
+_CCCL_AGGREGATE_ALL_OF_CASE(13);
+_CCCL_AGGREGATE_ALL_OF_CASE(14);
+_CCCL_AGGREGATE_ALL_OF_CASE(15);
+
+#undef _CCCL_AGGREGATE_ALL_OF_CASE
+#undef _CCCL_AGGR_PROBE
+
+// return true if
+// - T is an aggregate
+// - T has a number of members between 0 and __aggregate_max_arity
+// - T is not empty
+template <template <typename> class _Predicate,
+          typename _Tp,
+          bool = is_aggregate_v<_Tp> && (__aggregate_arity_v<_Tp> >= 0)
+              && (__aggregate_arity_v<_Tp> <= __aggregate_max_arity)
+              && ((__aggregate_arity_v<_Tp> != 0) || is_empty_v<_Tp>)>
+struct __aggregate_all_of : false_type
+{};
+
+// Applies a Predicate to every member reachable by aggregate initialization
+template <template <typename> class _Predicate, typename _Tp>
+struct __aggregate_all_of<_Predicate, _Tp, true>
+    : decltype(__aggregate_all_of_fn<__aggregate_arity_v<_Tp>>::template __call<_Predicate, _Tp>(0))
+{};
+
+_CCCL_END_NAMESPACE_CUDA_STD
+
+_CCCL_DIAG_POP
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // _CUDA_STD___TYPE_TRAITS_AGGREGATE_MEMBERS_ALL_OF_H
diff --git a/libcudacxx/include/cuda/std/string_view b/libcudacxx/include/cuda/std/string_view
index 2b8348b341f..74bdcee6400 100644
--- a/libcudacxx/include/cuda/std/string_view
+++ b/libcudacxx/include/cuda/std/string_view
@@ -51,8 +51,8 @@
 #include <cuda/std/__type_traits/is_convertible.h>
 #include <cuda/std/__type_traits/is_same.h>
 #include <cuda/std/__type_traits/is_standard_layout.h>
-#include <cuda/std/__type_traits/is_trivially_constructible.h>
 #include <cuda/std/__type_traits/is_trivially_copyable.h>
+#include <cuda/std/__type_traits/is_trivially_default_constructible.h>
 #include <cuda/std/__type_traits/remove_cvref.h>
 #include <cuda/std/__type_traits/type_identity.h>
 #include <cuda/std/__utility/ctad_support.h>
diff --git a/libcudacxx/include/cuda/type_traits b/libcudacxx/include/cuda/type_traits
index 446359b6700..f063eb0dafd 100644
--- a/libcudacxx/include/cuda/type_traits
+++ b/libcudacxx/include/cuda/type_traits
@@ -22,6 +22,7 @@
 #endif // no system header
 
 #include <cuda/__type_traits/is_floating_point.h>
+#include <cuda/__type_traits/is_trivially_copyable.h>
 #include <cuda/__type_traits/is_vector_type.h>
 #include <cuda/__type_traits/vector_type.h>
 #include <cuda/std/type_traits>
diff --git a/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.aggr.pass.cpp b/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.aggr.pass.cpp
new file mode 100644
index 00000000000..f668a29d2f5
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.aggr.pass.cpp
@@ -0,0 +1,114 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/type_traits>
+
+template <class _Tp>
+__host__ __device__ void test_is_trivially_copyable()
+{
+  static_assert(cuda::is_trivially_copyable<_Tp>::value);
+  static_assert(cuda::is_trivially_copyable<const _Tp>::value);
+  static_assert(cuda::is_trivially_copyable_v<_Tp>);
+  static_assert(cuda::is_trivially_copyable_v<const _Tp>);
+}
+
+template <class _Tp>
+__host__ __device__ void test_is_not_trivially_copyable()
+{
+  static_assert(!cuda::is_trivially_copyable<_Tp>::value);
+  static_assert(!cuda::is_trivially_copyable<const _Tp>::value);
+  static_assert(!cuda::is_trivially_copyable_v<_Tp>);
+  static_assert(!cuda::is_trivially_copyable_v<const _Tp>);
+}
+
+struct SingleMember
+{
+  int x;
+};
+
+struct DerivedStruct : SingleMember
+{
+  int y;
+};
+
+struct NesterStruct
+{
+  SingleMember x;
+};
+
+struct ArrayMember
+{
+  int values[2];
+};
+
+//----------------------------------------------------------------------------------------------------------------------
+// custom type
+
+struct CustomNonTrivialType
+{
+  int x;
+
+  CustomNonTrivialType() = default;
+  __host__ __device__ CustomNonTrivialType(const CustomNonTrivialType&) {}
+};
+
+template <>
+constexpr bool cuda::is_trivially_copyable_v<CustomNonTrivialType> = true;
+
+struct SingleMemberCustom
+{
+  CustomNonTrivialType x;
+};
+
+struct DerivedStructCustom : SingleMemberCustom
+{
+  CustomNonTrivialType y;
+};
+
+struct NestedStructCustom
+{
+  CustomNonTrivialType z;
+};
+
+struct ArrayMemberCustom
+{
+  CustomNonTrivialType values[2];
+};
+
+//----------------------------------------------------------------------------------------------------------------------
+// non trivially copyable type
+
+struct NonTriviallyCopyable
+{
+  __host__ __device__ NonTriviallyCopyable(const NonTriviallyCopyable&) {};
+};
+
+struct RelaxedWithNonRelaxedMember
+{
+  CustomNonTrivialType x;
+  NonTriviallyCopyable y;
+};
+
+__host__ __device__ void test()
+{
+  test_is_trivially_copyable<SingleMember>();
+  test_is_trivially_copyable<DerivedStruct>();
+  test_is_trivially_copyable<NesterStruct>();
+  test_is_trivially_copyable<ArrayMember>();
+  test_is_trivially_copyable<ArrayMember>();
+  test_is_trivially_copyable<DerivedStructCustom>();
+  test_is_trivially_copyable<NestedStructCustom>();
+  test_is_not_trivially_copyable<RelaxedWithNonRelaxedMember>();
+}
+
+int main(int, char**)
+{
+  test();
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.mem.pass.cpp b/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.mem.pass.cpp
new file mode 100644
index 00000000000..b2ed55e0979
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.mem.pass.cpp
@@ -0,0 +1,323 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda_bf16.h>
+#include <cuda_fp16.h>
+
+#include <cuda/__complex_>
+#include <cuda/std/array>
+#include <cuda/std/cassert>
+#include <cuda/std/complex>
+#include <cuda/std/cstring>
+#include <cuda/std/tuple>
+#include <cuda/std/utility>
+#include <cuda/type_traits>
+
+#include "test_macros.h"
+
+_CCCL_DIAG_PUSH
+_CCCL_DIAG_SUPPRESS_CLANG("-Wunused-local-typedef")
+
+// operator== for CUDA vector types and dim3 (not provided by the toolkit)
+template <class T, cuda::std::enable_if_t<cuda::is_vector_type_v<T>, int> = 0>
+__host__ __device__ bool operator==(T a, T b)
+{
+  if constexpr (sizeof(T) == sizeof(decltype(T::x)))
+  {
+    return a.x == b.x;
+  }
+  else if constexpr (sizeof(T) == 2 * sizeof(decltype(T::x)))
+  {
+    return a.x == b.x && a.y == b.y;
+  }
+  else if constexpr (sizeof(T) == 3 * sizeof(decltype(T::x)))
+  {
+    return a.x == b.x && a.y == b.y && a.z == b.z;
+  }
+  else
+  {
+    return a.x == b.x && a.y == b.y && a.z == b.z && a.w == b.w;
+  }
+}
+
+template <class T>
+__host__ __device__ bool operator==(cuda::complex<T> a, cuda::complex<T> b)
+{
+  return a.real() == b.real() && a.imag() == b.imag();
+}
+
+template <typename T>
+__host__ __device__ void test_memcpy_roundtrip(T from)
+{
+  static_assert(cuda::is_trivially_copyable_v<T>);
+  T to;
+  ::memcpy(static_cast<void*>(&to), static_cast<const void*>(&from), sizeof(T));
+  assert(from == to);
+}
+
+#define CAST(base_type, val) static_cast<decltype(base_type##1 ::x)>(val)
+
+#define REPEAT_1(base_type, index) CAST(base_type, input[index][0])
+#define REPEAT_2(base_type, index) REPEAT_1(base_type, index), CAST(base_type, input[index][1])
+#define REPEAT_3(base_type, index) REPEAT_2(base_type, index), CAST(base_type, input[index][2])
+#define REPEAT_4(base_type, index) REPEAT_3(base_type, index), CAST(base_type, input[index][3])
+
+#define TEST_CUDA_VECTOR_TYPE(base_type, size)           \
+  {                                                      \
+    for (base_type##size i :                             \
+         {base_type##size{REPEAT_##size(base_type, 0)},  \
+          base_type##size{REPEAT_##size(base_type, 1)},  \
+          base_type##size{REPEAT_##size(base_type, 2)},  \
+          base_type##size{REPEAT_##size(base_type, 3)},  \
+          base_type##size{REPEAT_##size(base_type, 4)},  \
+          base_type##size{REPEAT_##size(base_type, 5)},  \
+          base_type##size{REPEAT_##size(base_type, 6)}}) \
+    {                                                    \
+      test_memcpy_roundtrip(i);                          \
+    }                                                    \
+  }
+
+#define TEST_CUDA_VECTOR_TYPES(base_type) \
+  TEST_CUDA_VECTOR_TYPE(base_type, 1)     \
+  TEST_CUDA_VECTOR_TYPE(base_type, 2)     \
+  TEST_CUDA_VECTOR_TYPE(base_type, 3)     \
+  TEST_CUDA_VECTOR_TYPE(base_type, 4)
+
+__host__ __device__ bool tests()
+{
+  // standard scalar types
+  test_memcpy_roundtrip(42);
+  test_memcpy_roundtrip(0.0f);
+  test_memcpy_roundtrip(3.14159);
+  test_memcpy_roundtrip(static_cast<short>(7));
+  test_memcpy_roundtrip(static_cast<char>('A'));
+
+  // cuda::std::pair
+  using pair = cuda::std::pair<float, int>;
+  for (pair i :
+       {pair{0.0f, 1},
+        pair{1.0f, 2},
+        pair{-1.0f, 3},
+        pair{10.0f, 4},
+        pair{-10.0f, 5},
+        pair{2.71828f, 6},
+        pair{3.14159f, 7}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  // cuda::std::tuple
+  using tuple = cuda::std::tuple<int, float>;
+  for (tuple i :
+       {tuple{1, 0.0f},
+        tuple{2, 1.0f},
+        tuple{3, -1.0f},
+        tuple{4, 10.0f},
+        tuple{5, -10.0f},
+        tuple{6, 2.71828f},
+        tuple{7, 3.14159f}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  // cuda::std::array
+  using array = cuda::std::array<float, 2>;
+  for (array i :
+       {array{0.0f, 1.0f},
+        array{1.0f, 2.0f},
+        array{-1.0f, 3.0f},
+        array{10.0f, 4.0f},
+        array{-10.0f, 5.0f},
+        array{2.71828f, 6.0f},
+        array{3.14159f, 7.0f}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  // cuda::std::complex
+  using std_complex_f = cuda::std::complex<float>;
+  for (std_complex_f i :
+       {std_complex_f{0.0f, 1.0f},
+        std_complex_f{1.0f, -1.0f},
+        std_complex_f{-1.0f, 0.0f},
+        std_complex_f{10.0f, -10.0f},
+        std_complex_f{-10.0f, 10.0f},
+        std_complex_f{2.71828f, 3.14159f},
+        std_complex_f{3.14159f, 2.71828f}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  // cuda::complex
+  using cuda_complex_f = cuda::complex<float>;
+  for (cuda_complex_f i :
+       {cuda_complex_f{0.0f, 1.0f},
+        cuda_complex_f{1.0f, -1.0f},
+        cuda_complex_f{-1.0f, 0.0f},
+        cuda_complex_f{10.0f, -10.0f},
+        cuda_complex_f{-10.0f, 10.0f},
+        cuda_complex_f{2.71828f, 3.14159f},
+        cuda_complex_f{3.14159f, 2.71828f}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  // CUDA vector types
+  constexpr double input[7][4] = {
+    {0.0, 1.0, -7.0, -0.0},
+    {1.0, 2.0, -7.0, -1.0},
+    {-1.0, 3.0, -7.0, 1.0},
+    {10.0, 4.0, -7.0, -10.0},
+    {-10.0, 5.0, -7.0, 10.0},
+    {2.71828, 6.0, -7.0, -2.71828},
+    {3.14159, 7.0, -7.0, -3.14159}};
+
+  TEST_CUDA_VECTOR_TYPES(char)
+  TEST_CUDA_VECTOR_TYPES(short)
+  TEST_CUDA_VECTOR_TYPES(int)
+  TEST_CUDA_VECTOR_TYPES(float)
+
+#if !_CCCL_CUDA_COMPILER(CLANG)
+  using uchar  = unsigned char;
+  using ushort = unsigned short;
+  using uint   = unsigned int;
+  using ulong  = unsigned long;
+  TEST_CUDA_VECTOR_TYPES(uchar)
+  TEST_CUDA_VECTOR_TYPES(ushort)
+  TEST_CUDA_VECTOR_TYPES(uint)
+  TEST_CUDA_VECTOR_TYPE(ulong, 1)
+  TEST_CUDA_VECTOR_TYPE(ulong, 2)
+  TEST_CUDA_VECTOR_TYPE(ulong, 3)
+#endif // !_CCCL_CUDA_COMPILER(CLANG)
+
+  TEST_CUDA_VECTOR_TYPE(long, 1)
+  TEST_CUDA_VECTOR_TYPE(long, 2)
+  TEST_CUDA_VECTOR_TYPE(long, 3)
+
+  using longlong  = long long;
+  using ulonglong = unsigned long long;
+  TEST_CUDA_VECTOR_TYPE(longlong, 1)
+  TEST_CUDA_VECTOR_TYPE(longlong, 2)
+  TEST_CUDA_VECTOR_TYPE(longlong, 3)
+  TEST_CUDA_VECTOR_TYPE(ulonglong, 1)
+  TEST_CUDA_VECTOR_TYPE(ulonglong, 2)
+  TEST_CUDA_VECTOR_TYPE(ulonglong, 3)
+  TEST_CUDA_VECTOR_TYPE(double, 1)
+  TEST_CUDA_VECTOR_TYPE(double, 2)
+  TEST_CUDA_VECTOR_TYPE(double, 3)
+
+  for (dim3 i :
+       {dim3{0u, 1u, 2u},
+        dim3{1u, 2u, 3u},
+        dim3{10u, 20u, 30u},
+        dim3{100u, 200u, 300u},
+        dim3{255u, 128u, 64u},
+        dim3{1024u, 512u, 256u},
+        dim3{4096u, 2048u, 1024u}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  return true;
+}
+
+__host__ __device__ bool tests_nvfp()
+{
+#if _LIBCUDACXX_HAS_NVFP16()
+  for (__half i :
+       {__float2half(0.0f),
+        __float2half(1.0f),
+        __float2half(-1.0f),
+        __float2half(10.0f),
+        __float2half(-10.0f),
+        __float2half(2.71828f),
+        __float2half(3.14159f)})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  for (__half2 i :
+       {__half2{__float2half(0.0f), __float2half(1.0f)},
+        __half2{__float2half(-1.0f), __float2half(2.0f)},
+        __half2{__float2half(10.0f), __float2half(-10.0f)},
+        __half2{__float2half(2.71828f), __float2half(3.14159f)}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  test_memcpy_roundtrip(
+    cuda::std::array<__half, 4>{__float2half(1.0f), __float2half(2.0f), __float2half(3.0f), __float2half(4.0f)});
+  test_memcpy_roundtrip(cuda::std::pair<__half, __half>{__float2half(1.0f), __float2half(2.0f)});
+  test_memcpy_roundtrip(cuda::std::tuple<__half, __half>{__float2half(1.0f), __float2half(2.0f)});
+
+  test_memcpy_roundtrip(cuda::std::array<cuda::std::pair<__half, __half>, 2>{
+    cuda::std::pair<__half, __half>{__float2half(1.0f), __float2half(2.0f)},
+    cuda::std::pair<__half, __half>{__float2half(3.0f), __float2half(4.0f)}});
+
+  using complex_half = cuda::std::complex<__half>;
+  for (complex_half i :
+       {complex_half{__float2half(0.0f), __float2half(1.0f)},
+        complex_half{__float2half(1.0f), __float2half(-1.0f)},
+        complex_half{__float2half(-1.0f), __float2half(0.0f)},
+        complex_half{__float2half(10.0f), __float2half(-10.0f)}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+#endif // _LIBCUDACXX_HAS_NVFP16()
+
+#if _LIBCUDACXX_HAS_NVBF16()
+  for (__nv_bfloat16 i :
+       {__float2bfloat16(0.0f),
+        __float2bfloat16(1.0f),
+        __float2bfloat16(-1.0f),
+        __float2bfloat16(10.0f),
+        __float2bfloat16(-10.0f),
+        __float2bfloat16(2.71828f),
+        __float2bfloat16(3.14159f)})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  for (__nv_bfloat162 i :
+       {__nv_bfloat162{__float2bfloat16(0.0f), __float2bfloat16(1.0f)},
+        __nv_bfloat162{__float2bfloat16(-1.0f), __float2bfloat16(2.0f)},
+        __nv_bfloat162{__float2bfloat16(10.0f), __float2bfloat16(-10.0f)},
+        __nv_bfloat162{__float2bfloat16(2.71828f), __float2bfloat16(3.14159f)}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+
+  test_memcpy_roundtrip(cuda::std::array<__nv_bfloat16, 2>{__float2bfloat16(1.0f), __float2bfloat16(2.0f)});
+  test_memcpy_roundtrip(cuda::std::pair<__nv_bfloat16, __nv_bfloat16>{__float2bfloat16(1.0f), __float2bfloat16(2.0f)});
+
+  using complex_bfloat = cuda::std::complex<__nv_bfloat16>;
+  for (complex_bfloat i :
+       {complex_bfloat{__float2bfloat16(0.0f), __float2bfloat16(1.0f)},
+        complex_bfloat{__float2bfloat16(1.0f), __float2bfloat16(-1.0f)},
+        complex_bfloat{__float2bfloat16(-1.0f), __float2bfloat16(0.0f)},
+        complex_bfloat{__float2bfloat16(10.0f), __float2bfloat16(-10.0f)}})
+  {
+    test_memcpy_roundtrip(i);
+  }
+#endif // _LIBCUDACXX_HAS_NVBF16()
+
+#if _LIBCUDACXX_HAS_NVFP16() && _LIBCUDACXX_HAS_NVBF16()
+  test_memcpy_roundtrip(cuda::std::tuple<__half, __nv_bfloat16>{__float2half(1.0f), __float2bfloat16(2.0f)});
+#endif // _LIBCUDACXX_HAS_NVFP16() && _LIBCUDACXX_HAS_NVBF16()
+  return true;
+}
+
+int main(int, char**)
+{
+  assert(tests());
+  assert(tests_nvfp());
+  return 0;
+}
+
+_CCCL_DIAG_POP
diff --git a/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.pass.cpp b/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.pass.cpp
new file mode 100644
index 00000000000..bb0309fa96e
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.pass.cpp
@@ -0,0 +1,142 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda_bf16.h>
+#include <cuda_fp16.h>
+
+#include <cuda/__complex_>
+#include <cuda/std/array>
+#include <cuda/std/complex>
+#include <cuda/std/tuple>
+#include <cuda/std/utility>
+#include <cuda/type_traits>
+
+#include "test_macros.h"
+
+template <class T>
+__host__ __device__ void test_is_trivially_copyable()
+{
+  static_assert(cuda::is_trivially_copyable<T>::value);
+  static_assert(cuda::is_trivially_copyable<const T>::value);
+  static_assert(cuda::is_trivially_copyable_v<T>);
+  static_assert(cuda::is_trivially_copyable_v<const T>);
+}
+
+template <class T>
+__host__ __device__ void test_is_not_trivially_copyable()
+{
+  static_assert(!cuda::is_trivially_copyable<T>::value);
+  static_assert(!cuda::is_trivially_copyable<const T>::value);
+  static_assert(!cuda::is_trivially_copyable_v<T>);
+  static_assert(!cuda::is_trivially_copyable_v<const T>);
+}
+
+struct TrivialPod
+{
+  int x;
+  float y;
+};
+
+class NonTriviallyCopyable
+{
+public:
+  __host__ __device__ NonTriviallyCopyable(const NonTriviallyCopyable&) {}
+};
+
+__host__ __device__ void test()
+{
+  // standard trivially copyable types
+  test_is_trivially_copyable<int>();
+  test_is_trivially_copyable<float>();
+  test_is_trivially_copyable<double>();
+  test_is_trivially_copyable<TrivialPod>();
+
+  // C-style arrays of trivially copyable types
+  static_assert(cuda::is_trivially_copyable_v<int[4]>);
+  static_assert(cuda::is_trivially_copyable_v<const int[4]>);
+  static_assert(cuda::is_trivially_copyable_v<TrivialPod[2]>);
+
+  // cuda::std::array, pair, tuple of trivially copyable types
+  test_is_trivially_copyable<cuda::std::array<int, 4>>();
+  test_is_trivially_copyable<cuda::std::pair<int, float>>();
+  test_is_trivially_copyable<cuda::std::tuple<int, float>>();
+  test_is_trivially_copyable<cuda::std::tuple<>>();
+
+  // complex types
+  test_is_trivially_copyable<cuda::std::complex<float>>();
+  test_is_trivially_copyable<cuda::complex<float>>();
+
+#if _CCCL_HAS_NVFP16()
+  test_is_trivially_copyable<cuda::std::complex<__half>>();
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVBF16()
+  test_is_trivially_copyable<cuda::std::complex<__nv_bfloat16>>();
+#endif // _CCCL_HAS_NVBF16()
+
+  // extended floating point scalar types
+#if _CCCL_HAS_NVFP16()
+  test_is_trivially_copyable<__half>();
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVBF16()
+  test_is_trivially_copyable<__nv_bfloat16>();
+#endif // _CCCL_HAS_NVBF16()
+#if _CCCL_HAS_NVFP8_E4M3()
+  test_is_trivially_copyable<__nv_fp8_e4m3>();
+#endif // _CCCL_HAS_NVFP8_E4M3()
+
+  // extended floating point vector types
+#if _CCCL_HAS_NVFP16()
+  test_is_trivially_copyable<__half2>();
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVBF16()
+  test_is_trivially_copyable<__nv_bfloat162>();
+#endif // _CCCL_HAS_NVBF16()
+#if _CCCL_HAS_NVFP8()
+  test_is_trivially_copyable<__nv_fp8x2_e4m3>();
+#endif // _CCCL_HAS_NVFP8()
+
+  // padding-free compositions of extended floating point types
+#if _CCCL_HAS_NVFP16()
+  static_assert(cuda::is_trivially_copyable_v<__half[4]>);
+  static_assert(cuda::is_trivially_copyable_v<const __half[4]>);
+  test_is_trivially_copyable<cuda::std::array<__half, 4>>();
+  test_is_trivially_copyable<cuda::std::pair<__half, __half>>();
+  test_is_trivially_copyable<cuda::std::tuple<__half, __half>>();
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVBF16()
+  test_is_trivially_copyable<cuda::std::array<__nv_bfloat16, 2>>();
+  test_is_trivially_copyable<cuda::std::pair<__nv_bfloat16, __nv_bfloat16>>();
+#endif // _CCCL_HAS_NVBF16()
+
+  // compositions with padding
+#if _CCCL_HAS_NVFP16()
+  test_is_trivially_copyable<cuda::std::pair<__half, int>>();
+  test_is_trivially_copyable<cuda::std::pair<__half, float>>();
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVBF16()
+  test_is_trivially_copyable<cuda::std::pair<__nv_bfloat16, int>>();
+#endif // _CCCL_HAS_NVBF16()
+
+  // nested compositions
+#if _CCCL_HAS_NVFP16()
+  test_is_trivially_copyable<cuda::std::array<cuda::std::pair<__half, __half>, 2>>();
+#endif // _CCCL_HAS_NVFP16()
+#if _CCCL_HAS_NVFP16() && _CCCL_HAS_NVBF16()
+  test_is_trivially_copyable<cuda::std::tuple<__half, __nv_bfloat16>>();
+#endif // _CCCL_HAS_NVFP16() && _CCCL_HAS_NVBF16()
+
+  // non-trivially copyable types
+  test_is_not_trivially_copyable<NonTriviallyCopyable>();
+}
+
+int main(int, char**)
+{
+  test();
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/bit/bit.cast/bit_cast.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.cast/bit_cast.pass.cpp
index 4c9af27e726..4b165acdefa 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/bit/bit.cast/bit_cast.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.cast/bit_cast.pass.cpp
@@ -17,9 +17,12 @@
 #include <cuda/std/bit>
 #include <cuda/std/cassert>
 #include <cuda/std/cmath>
+#include <cuda/std/complex>
 #include <cuda/std/cstdint>
 #include <cuda/std/cstring>
 #include <cuda/std/limits>
+#include <cuda/std/tuple>
+#include <cuda/std/utility>
 
 #include "test_macros.h"
 
@@ -53,7 +56,7 @@ __host__ __device__ void test_roundtrip_through_nested_T(T from)
   {
     T x;
   };
-  static_assert(sizeof(Nested) == sizeof(T), "");
+  static_assert(sizeof(Nested) == sizeof(T));
 
   Nested middle  = cuda::std::bit_cast<Nested>(from);
   T to           = cuda::std::bit_cast<T>(middle);
@@ -72,7 +75,7 @@ __host__ __device__ void test_roundtrip_through_nested_T(T from)
 template <typename Intermediate, bool HasUniqueObjectRepresentations = true, typename T>
 __host__ __device__ void test_roundtrip_through(T from)
 {
-  static_assert(sizeof(Intermediate) == sizeof(T), "");
+  static_assert(sizeof(Intermediate) == sizeof(T));
 
   Intermediate middle  = cuda::std::bit_cast<Intermediate>(from);
   T to                 = cuda::std::bit_cast<T>(middle);
@@ -88,6 +91,16 @@ __host__ __device__ void test_roundtrip_through(T from)
   }
 }
 
+struct TrivialPod
+{
+  int x;
+  float y;
+  __host__ __device__ friend bool operator==(TrivialPod a, TrivialPod b)
+  {
+    return a.x == b.x && a.y == b.y;
+  }
+};
+
 template <typename T>
 __host__ __device__ _CCCL_CONSTEXPR_BIT_CAST cuda::std::array<T, 10> generate_signed_integral_values()
 {
@@ -319,6 +332,179 @@ __host__ __device__ bool tests()
   }
 #endif // _LIBCUDACXX_HAS_NVBF16()
 
+  // User-defined trivially copyable type
+  for (TrivialPod i : {TrivialPod{0, 0.0f}, TrivialPod{1, 1.0f}, TrivialPod{-1, 3.14159f}, TrivialPod{42, 2.71828f}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+
+  // cuda::std::array
+  for (cuda::std::array<int, 4> i :
+       {cuda::std::array<int, 4>{0, 1, 2, 3},
+        cuda::std::array<int, 4>{-1, -2, -3, -4},
+        cuda::std::array<int, 4>{100, 200, 300, 400}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+
+  // cuda::std::pair
+  for (cuda::std::pair<int, float> i :
+       {cuda::std::pair<int, float>{0, 0.0f},
+        cuda::std::pair<int, float>{1, 1.0f},
+        cuda::std::pair<int, float>{-1, 3.14159f}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+
+  // cuda::std::tuple
+  for (cuda::std::tuple<int, float> i :
+       {cuda::std::tuple<int, float>{0, 0.0f},
+        cuda::std::tuple<int, float>{1, 1.0f},
+        cuda::std::tuple<int, float>{-1, 3.14159f}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+
+  // cuda::std::tuple<> (empty, sizeof == 1 with no data bytes)
+  test_roundtrip_through_buffer<false>(cuda::std::tuple<>{});
+
+  // cuda::std::complex<float>
+  for (cuda::std::complex<float> i :
+       {cuda::std::complex<float>{0.0f, 1.0f},
+        cuda::std::complex<float>{1.0f, -1.0f},
+        cuda::std::complex<float>{-1.0f, 0.0f},
+        cuda::std::complex<float>{10.0f, -10.0f},
+        cuda::std::complex<float>{2.71828f, 3.14159f}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+
+#if _LIBCUDACXX_HAS_NVFP16()
+  // cuda::std::complex<__half>
+  for (cuda::std::complex<__half> i :
+       {cuda::std::complex<__half>{__float2half(0.0f), __float2half(1.0f)},
+        cuda::std::complex<__half>{__float2half(1.0f), __float2half(-1.0f)},
+        cuda::std::complex<__half>{__float2half(-1.0f), __float2half(0.0f)},
+        cuda::std::complex<__half>{__float2half(10.0f), __float2half(-10.0f)}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+#endif // _LIBCUDACXX_HAS_NVFP16()
+
+#if _LIBCUDACXX_HAS_NVBF16()
+  // cuda::std::complex<__nv_bfloat16>
+  for (cuda::std::complex<__nv_bfloat16> i :
+       {cuda::std::complex<__nv_bfloat16>{__float2bfloat16(0.0f), __float2bfloat16(1.0f)},
+        cuda::std::complex<__nv_bfloat16>{__float2bfloat16(1.0f), __float2bfloat16(-1.0f)},
+        cuda::std::complex<__nv_bfloat16>{__float2bfloat16(-1.0f), __float2bfloat16(0.0f)},
+        cuda::std::complex<__nv_bfloat16>{__float2bfloat16(10.0f), __float2bfloat16(-10.0f)}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+#endif // _LIBCUDACXX_HAS_NVBF16()
+
+  // Extended floating point vector types
+#if _LIBCUDACXX_HAS_NVFP16()
+  for (__half2 i :
+       {__half2{__float2half(0.0f), __float2half(1.0f)},
+        __half2{__float2half(-1.0f), __float2half(2.0f)},
+        __half2{__float2half(10.0f), __float2half(-10.0f)},
+        __half2{__float2half(2.71828f), __float2half(3.14159f)}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+#endif // _LIBCUDACXX_HAS_NVFP16()
+
+#if _LIBCUDACXX_HAS_NVBF16()
+  for (__nv_bfloat162 i :
+       {__nv_bfloat162{__float2bfloat16(0.0f), __float2bfloat16(1.0f)},
+        __nv_bfloat162{__float2bfloat16(-1.0f), __float2bfloat16(2.0f)},
+        __nv_bfloat162{__float2bfloat16(10.0f), __float2bfloat16(-10.0f)},
+        __nv_bfloat162{__float2bfloat16(2.71828f), __float2bfloat16(3.14159f)}})
+  {
+    test_roundtrip_through_nested_T(i);
+    test_roundtrip_through_buffer(i);
+  }
+#endif // _LIBCUDACXX_HAS_NVBF16()
+
+  // Padding-free compositions of extended floating point scalar types
+#if _LIBCUDACXX_HAS_NVFP16()
+  {
+    const auto arr =
+      cuda::std::array<__half, 4>{__float2half(1.0f), __float2half(2.0f), __float2half(3.0f), __float2half(4.0f)};
+    test_roundtrip_through_nested_T(arr);
+    test_roundtrip_through_buffer(arr);
+  }
+  {
+    const auto p = cuda::std::pair<__half, __half>{__float2half(1.0f), __float2half(2.0f)};
+    test_roundtrip_through_nested_T(p);
+    test_roundtrip_through_buffer(p);
+  }
+  {
+    const auto t = cuda::std::tuple<__half, __half>{__float2half(1.0f), __float2half(2.0f)};
+    test_roundtrip_through_nested_T(t);
+    test_roundtrip_through_buffer(t);
+  }
+  {
+    const auto nested = cuda::std::array<cuda::std::pair<__half, __half>, 2>{
+      cuda::std::pair<__half, __half>{__float2half(1.0f), __float2half(2.0f)},
+      cuda::std::pair<__half, __half>{__float2half(3.0f), __float2half(4.0f)}};
+    test_roundtrip_through_nested_T(nested);
+    test_roundtrip_through_buffer(nested);
+  }
+#endif // _LIBCUDACXX_HAS_NVFP16()
+
+#if _LIBCUDACXX_HAS_NVBF16()
+  {
+    const auto arr = cuda::std::array<__nv_bfloat16, 2>{__float2bfloat16(1.0f), __float2bfloat16(2.0f)};
+    test_roundtrip_through_nested_T(arr);
+    test_roundtrip_through_buffer(arr);
+  }
+  {
+    const auto p = cuda::std::pair<__nv_bfloat16, __nv_bfloat16>{__float2bfloat16(1.0f), __float2bfloat16(2.0f)};
+    test_roundtrip_through_nested_T(p);
+    test_roundtrip_through_buffer(p);
+  }
+#endif // _LIBCUDACXX_HAS_NVBF16()
+
+#if _LIBCUDACXX_HAS_NVFP16() && _LIBCUDACXX_HAS_NVBF16()
+  {
+    const auto t = cuda::std::tuple<__half, __nv_bfloat16>{__float2half(1.0f), __float2bfloat16(2.0f)};
+    test_roundtrip_through_nested_T(t);
+    test_roundtrip_through_buffer(t);
+  }
+#endif // _LIBCUDACXX_HAS_NVFP16() && _LIBCUDACXX_HAS_NVBF16()
+
+  // Padded compositions
+#if _LIBCUDACXX_HAS_NVFP16()
+  {
+    const auto p1 = cuda::std::pair<__half, int>{__float2half(1.0f), 42};
+    test_roundtrip_through_nested_T<false>(p1);
+    test_roundtrip_through_buffer<false>(p1);
+  }
+  {
+    const auto p2 = cuda::std::pair<__half, float>{__float2half(1.0f), 3.14f};
+    test_roundtrip_through_nested_T<false>(p2);
+    test_roundtrip_through_buffer<false>(p2);
+  }
+#endif // _LIBCUDACXX_HAS_NVFP16()
+
+#if _LIBCUDACXX_HAS_NVBF16()
+  {
+    const auto p = cuda::std::pair<__nv_bfloat16, int>{__float2bfloat16(1.0f), 42};
+    test_roundtrip_through_nested_T<false>(p);
+    test_roundtrip_through_buffer<false>(p);
+  }
+#endif // _LIBCUDACXX_HAS_NVBF16()
+
   // Test pointers
   {
     {
@@ -362,7 +548,7 @@ int main(int, char**)
 {
   tests();
 #if defined(_CCCL_BUILTIN_BIT_CAST)
-  static_assert(basic_constexpr_test(), "");
+  static_assert(basic_constexpr_test());
 #endif // _CCCL_BUILTIN_BIT_CAST
   return 0;
 }
diff --git a/libcudacxx/test/libcudacxx/std/utilities/meta/meta.unary/meta.unary.prop/aggregate_members_all_of.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/meta/meta.unary/meta.unary.prop/aggregate_members_all_of.pass.cpp
new file mode 100644
index 00000000000..30274fa6525
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/utilities/meta/meta.unary/meta.unary.prop/aggregate_members_all_of.pass.cpp
@@ -0,0 +1,201 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/__type_traits/aggregate_members_all_of.h>
+#include <cuda/std/__type_traits/is_integral.h>
+#include <cuda/std/__type_traits/is_trivially_copyable.h>
+
+#include "test_macros.h"
+
+//----------------------------------------------------------------------------------------------------------------------
+// test types
+
+struct Empty
+{};
+
+struct OneMember
+{
+  int x;
+};
+
+struct TwoMembers
+{
+  int x;
+  float y;
+};
+
+struct ThreeMembers
+{
+  int x;
+  float y;
+  double z;
+};
+
+struct AllInts
+{
+  int a;
+  int b;
+  int c;
+};
+
+struct AllFloats
+{
+  float a;
+  double b;
+};
+
+struct MixedIntFloat
+{
+  int a;
+  float b;
+};
+
+struct LargeAggregate
+{
+  int m0;
+  int m1;
+  int m2;
+  int m3;
+  int m4;
+  int m5;
+  int m6;
+  int m7;
+  int m8;
+  int m9;
+  int m10;
+  int m11;
+  int m12;
+  int m13;
+  int m14;
+  int m15;
+};
+
+struct TooLarge
+{
+  int m0;
+  int m1;
+  int m2;
+  int m3;
+  int m4;
+  int m5;
+  int m6;
+  int m7;
+  int m8;
+  int m9;
+  int m10;
+  int m11;
+  int m12;
+  int m13;
+  int m14;
+  int m15;
+  int m16;
+};
+
+struct Nested
+{
+  OneMember inner;
+  int extra;
+};
+
+struct WithArray
+{
+  int values[4];
+};
+
+class NonAggregate
+{
+public:
+  __host__ __device__ NonAggregate() {}
+  int x;
+};
+
+struct NonTriviallyCopyable
+{
+  NonTriviallyCopyable() = default;
+  __host__ __device__ NonTriviallyCopyable(const NonTriviallyCopyable&) {}
+  int x;
+};
+
+struct HasNonTriviallyCopyableMember
+{
+  int x;
+  NonTriviallyCopyable y;
+};
+
+//----------------------------------------------------------------------------------------------------------------------
+// predicates
+
+template <typename _Tp>
+using is_integral_pred = cuda::std::is_integral<_Tp>;
+
+template <typename _Tp>
+struct always_true : cuda::std::true_type
+{};
+
+template <typename _Tp>
+struct always_false : cuda::std::false_type
+{};
+
+template <typename _Tp>
+using is_trivially_copyable_pred = cuda::std::is_trivially_copyable<_Tp>;
+
+//----------------------------------------------------------------------------------------------------------------------
+// __aggregate_arity_v tests
+
+static_assert(cuda::std::__aggregate_arity_v<Empty> == 0);
+static_assert(cuda::std::__aggregate_arity_v<OneMember> == 1);
+static_assert(cuda::std::__aggregate_arity_v<TwoMembers> == 2);
+static_assert(cuda::std::__aggregate_arity_v<ThreeMembers> == 3);
+static_assert(cuda::std::__aggregate_arity_v<AllInts> == 3);
+static_assert(cuda::std::__aggregate_arity_v<AllFloats> == 2);
+static_assert(cuda::std::__aggregate_arity_v<LargeAggregate> == 16);
+static_assert(cuda::std::__aggregate_arity_v<NonAggregate> == -1);
+static_assert(cuda::std::__aggregate_arity_v<Nested> == 2);
+static_assert(cuda::std::__aggregate_arity_v<WithArray> == 4);
+
+//----------------------------------------------------------------------------------------------------------------------
+// __aggregate_all_of tests
+
+// empty aggregate: always true (vacuously)
+static_assert(cuda::std::__aggregate_all_of<always_true, Empty>::value);
+static_assert(cuda::std::__aggregate_all_of<always_false, Empty>::value);
+
+// all members satisfy the predicate
+static_assert(cuda::std::__aggregate_all_of<is_integral_pred, AllInts>::value);
+static_assert(cuda::std::__aggregate_all_of<always_true, ThreeMembers>::value);
+
+// not all members satisfy the predicate
+static_assert(!cuda::std::__aggregate_all_of<is_integral_pred, MixedIntFloat>::value);
+static_assert(!cuda::std::__aggregate_all_of<always_false, OneMember>::value);
+static_assert(!cuda::std::__aggregate_all_of<is_integral_pred, AllFloats>::value);
+
+// max arity aggregate
+static_assert(cuda::std::__aggregate_all_of<is_integral_pred, LargeAggregate>::value);
+
+// too many members: returns false (arity exceeds __aggregate_max_arity)
+static_assert(!cuda::std::__aggregate_all_of<is_integral_pred, TooLarge>::value);
+
+// nested aggregate / array members: brace elision lets the predicate see the flat elements
+static_assert(cuda::std::__aggregate_all_of<is_integral_pred, Nested>::value);
+static_assert(cuda::std::__aggregate_all_of<is_integral_pred, WithArray>::value);
+static_assert(cuda::std::__aggregate_all_of<is_trivially_copyable_pred, Nested>::value);
+static_assert(cuda::std::__aggregate_all_of<is_trivially_copyable_pred, WithArray>::value);
+
+// non-aggregate: always false
+static_assert(!cuda::std::__aggregate_all_of<always_true, NonAggregate>::value);
+static_assert(!cuda::std::__aggregate_all_of<is_integral_pred, NonAggregate>::value);
+
+// aggregate with a non-trivially-copyable member
+static_assert(!cuda::std::__aggregate_all_of<is_trivially_copyable_pred, HasNonTriviallyCopyableMember>::value);
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/thrust/thrust/detail/allocator/value_initialize_range.h b/thrust/thrust/detail/allocator/value_initialize_range.h
index 7986bfc07fe..0e2ad4ca6e4 100644
--- a/thrust/thrust/detail/allocator/value_initialize_range.h
+++ b/thrust/thrust/detail/allocator/value_initialize_range.h
@@ -20,6 +20,7 @@
 #include <thrust/uninitialized_fill.h>
 
 #include <cuda/std/__memory/allocator_traits.h>
+#include <cuda/std/__type_traits/is_trivially_default_constructible.h>
 
 THRUST_NAMESPACE_BEGIN
 namespace detail