cuda::is_trivially_copyable

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

docs/libcudacxx/extended_api/type_traits/is_trivially_copyable.rst

@@ -0,0 +1,103 @@
+.. _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.
+- Aggregates: trivially copyable when all members are.
+
+``const`` qualification is handled transparently, while ``volatile`` is compiler dependent.
+
+Custom Specialization
+---------------------
+
+Users may specialize ``cuda::is_trivially_copyable_v`` for types whose semantics allow copying with ``memcpy``, but which the compiler does not consider to be trivially copyable.
+
+A `trivially copyable <https://en.cppreference.com/w/cpp/language/classes.html>`__ class is a class that
+
+- has at least one eligible copy constructor, move constructor, copy assignment operator, or move assignment operator,
+- each of its eligible copy constructors is trivial
+- each eligible move constructor is trivial
+- each eligible copy assignment operator is trivial
+- each eligible move assignment operator is trivial, and
+- has a non-deleted trivial destructor.
+
+.. warning::
+
+    The user is responsible for ensuring that the type is actually trivially copyable when specializing this variable template. Otherwise, the behavior is undefined.
+
+A common case is a user-declared type that has a copy constructor but is technically trivially copyable:
+
+.. code:: cuda
+
+    struct NonTriviallyCopyable {
+        int x;
+        __host__ __device__ NonTriviallyCopyable(const NonTriviallyCopyable&) {}
+    };
+
+    // Specializing the variable template
+    template <>
+    constexpr bool cuda::is_trivially_copyable_v<NonTriviallyCopyable> = true;
+
+    static_assert(cuda::is_trivially_copyable_v<NonTriviallyCopyable>);
+
+Examples
+--------
+
+.. code:: cuda
+
+   #include <cuda/type_traits>
+   #include <cuda/std/array>
+   #include <cuda/std/tuple>
+   #include <cuda/std/utility>
+
+   #include <cuda_fp16.h>
+
+   struct UserType {
+     __half x, y;
+   };
+
+   // 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>);
+
+   // Padding-free 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>>);
+
+   // Composites with padding are not trivially copyable
+   static_assert(!cuda::is_trivially_copyable_v<cuda::std::pair<__half, int>>);

libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h

@@ -0,0 +1,84 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+
+#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/__type_traits/is_vector_type.h>
+#include <cuda/std/__cstddef/types.h>
+#include <cuda/std/__fwd/array.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>
+constexpr bool __is_aggregate_trivially_copyable_v = false;
+
+//! Users are allowed to specialize this variable template for their own types
+template <typename _Tp>
+constexpr bool is_trivially_copyable_v =
+  ::cuda::std::is_trivially_copyable_v<::cuda::std::remove_const_t<_Tp>>
+  || ::cuda::std::__is_extended_floating_point_v<::cuda::std::remove_const_t<_Tp>>
+#if _CCCL_HAS_CTK()
+  || ::cuda::is_extended_fp_vector_type_v<::cuda::std::remove_const_t<_Tp>>
+#endif // _CCCL_HAS_CTK()
+  || __is_aggregate_trivially_copyable_v<::cuda::std::remove_const_t<_Tp>>;
+
+template <typename _Tp>
+constexpr bool is_trivially_copyable_v<_Tp[]> = is_trivially_copyable_v<_Tp>;
+
+template <typename _Tp, ::cuda::std::size_t _Size>
+constexpr bool is_trivially_copyable_v<_Tp[_Size]> = is_trivially_copyable_v<_Tp>;
+
+template <typename _Tp, ::cuda::std::size_t _Size>
+constexpr bool is_trivially_copyable_v<::cuda::std::array<_Tp, _Size>> = is_trivially_copyable_v<_Tp>;
+
+template <typename _T1, typename _T2>
+constexpr bool is_trivially_copyable_v<::cuda::std::pair<_T1, _T2>> =
+  is_trivially_copyable_v<_T1> && is_trivially_copyable_v<_T2>;
+
+template <typename... _Ts>
+constexpr bool is_trivially_copyable_v<::cuda::std::tuple<_Ts...>> = (is_trivially_copyable_v<_Ts> && ...);
+
+// 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>
+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;
+
+// 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

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>
+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>
+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

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>