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_relaxed
    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_relaxed <libcudacxx-extended-api-type_traits-is_trivially_copyable_relaxed>`
+     - 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_relaxed.rst

@@ -0,0 +1,110 @@
+.. _libcudacxx-extended-api-type_traits-is_trivially_copyable_relaxed:
+
+``cuda::is_trivially_copyable_relaxed``
+=======================================
+
+Defined in the ``<cuda/type_traits>`` header.
+
+.. code:: cuda
+
+   namespace cuda {
+
+   template <typename T>
+   constexpr bool is_trivially_copyable_relaxed_v = /* see below */;
+
+   template <typename T>
+   using is_trivially_copyable_relaxed = cuda::std::bool_constant<is_trivially_copyable_relaxed_v<T>>;
+
+   } // namespace cuda
+
+``cuda::is_trivially_copyable_relaxed_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_relaxed`` 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 relaxed trivially copyable when ``T`` is.
+- ``cuda::std::array<T, N>``: relaxed trivially copyable when ``T`` is.
+- ``cuda::std::pair<T1, T2>``: relaxed trivially copyable when both ``T1`` and ``T2`` are and the object has no padding.
+- ``cuda::std::tuple<Ts...>``: relaxed trivially copyable when all ``Ts...`` are and the object has no padding.
+
+``const`` qualification is handled transparently, while ``volatile`` is compiler dependent.
+
+.. note::
+
+  The type trait cannot determine if a structure (``struct`` or ``class``) contains extended floating-point types, and thus it recognizes the type as *trivially copyable*. The user must manually specialize the type trait for such types.
+
+Custom Specialization
+---------------------
+
+Users may specialize ``cuda::is_trivially_copyable_relaxed_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 eligible copy constructor 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 type that wraps extended floating-point fields and provides user-defined copy operations
+solely to add ``__host__ __device__`` annotations:
+
+.. code:: cuda
+
+    struct HalfWrapper {
+        __half value;
+    };
+
+    struct NonTriviallyCopyable {
+        __host__ __device__ NonTriviallyCopyable(const NonTriviallyCopyable&) {}
+    };
+
+    // Specializing the variable template
+    template <>
+    constexpr bool cuda::is_trivially_copyable_relaxed_v<HalfWrapper> = true;
+
+    template <>
+    constexpr bool cuda::is_trivially_copyable_relaxed_v<NonTriviallyCopyable> = true;
+
+    static_assert(cuda::is_trivially_copyable_relaxed_v<HalfWrapper>);
+    static_assert(cuda::is_trivially_copyable_relaxed_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>
+
+   // Standard trivially copyable types
+   static_assert(cuda::is_trivially_copyable_relaxed_v<int>);
+   static_assert(cuda::is_trivially_copyable_relaxed_v<float>);
+
+   // Extended floating-point types
+   static_assert(cuda::is_trivially_copyable_relaxed_v<__half>);
+   static_assert(cuda::is_trivially_copyable_relaxed_v<__nv_bfloat16>);
+   static_assert(cuda::is_trivially_copyable_relaxed_v<__half2>);
+
+   // Padding-free composite types containing extended floating-point types
+   static_assert(cuda::is_trivially_copyable_relaxed_v<__half[4]>);
+   static_assert(cuda::is_trivially_copyable_relaxed_v<cuda::std::array<__half, 4>>);
+   static_assert(cuda::is_trivially_copyable_relaxed_v<cuda::std::pair<__half, __half>>);
+   static_assert(cuda::is_trivially_copyable_relaxed_v<cuda::std::tuple<__half, __half>>);
+
+   // Composites with padding are not trivially copyable relaxed
+   static_assert(!cuda::is_trivially_copyable_relaxed_v<cuda::std::pair<__half, int>>);

libcudacxx/include/cuda/__type_traits/is_trivially_copyable_relaxed.h

@@ -0,0 +1,72 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CUDA__TYPE_TRAITS_IS_TRIVIALLY_COPYABLE_RELAXED_H
+#define __CUDA__TYPE_TRAITS_IS_TRIVIALLY_COPYABLE_RELAXED_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/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
+
+//! Users are allowed to specialize this variable template for their own types
+template <typename _Tp>
+constexpr bool is_trivially_copyable_relaxed_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()
+  ;
+
+template <typename _Tp>
+constexpr bool is_trivially_copyable_relaxed_v<_Tp[]> = is_trivially_copyable_relaxed_v<_Tp>;
+
+template <typename _Tp, ::cuda::std::size_t _Size>
+constexpr bool is_trivially_copyable_relaxed_v<_Tp[_Size]> = is_trivially_copyable_relaxed_v<_Tp>;
+
+template <typename _Tp, ::cuda::std::size_t _Size>
+constexpr bool is_trivially_copyable_relaxed_v<::cuda::std::array<_Tp, _Size>> = is_trivially_copyable_relaxed_v<_Tp>;
+
+template <typename _T1, typename _T2>
+constexpr bool is_trivially_copyable_relaxed_v<::cuda::std::pair<_T1, _T2>> =
+  is_trivially_copyable_relaxed_v<_T1> && is_trivially_copyable_relaxed_v<_T2>;
+
+template <typename... _Ts>
+constexpr bool is_trivially_copyable_relaxed_v<::cuda::std::tuple<_Ts...>> =
+  (is_trivially_copyable_relaxed_v<_Ts> && ...);
+
+// defined as alias so users cannot specialize it (they should specialize the variable template instead)
+template <typename _Tp>
+using is_trivially_copyable_relaxed = ::cuda::std::bool_constant<is_trivially_copyable_relaxed_v<_Tp>>;
+
+_CCCL_END_NAMESPACE_CUDA
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // __CUDA__TYPE_TRAITS_IS_TRIVIALLY_COPYABLE_RELAXED_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_relaxed.h>
 #include <cuda/__type_traits/is_vector_type.h>
 #include <cuda/__type_traits/vector_type.h>
 #include <cuda/std/type_traits>