`cuda::is_trivially_copyable` by fbusato · Pull Request #8265 · NVIDIA/cccl

fbusato · 2026-04-01T19:51:06Z

Description

Followup of the discussion in

Recover warp_shuffle original behavior (revert #8210) #8254

The PR introduces cuda::is_trivially_copyable_relaxed to support types that are actually trivially copyable but not recognized by the C++ std::is_trivially_copyable type trait.

The new trait supports:

trivially copyable types.
extended floating-point types.
extended floating-point vector types.
Raw arrays, cuda::std::array.
cuda::std::pair.
cuda::std::tuple.
Nested compositions of the above cases.

Potentially affected paths by `std::is_trivially_copyable`

unsupported cases:

cuda/std/__atomic/types/base.h:39
cuda/std/__atomic/types/reference.h:39
cuda/__memcpy_async/memcpy_async_barrier.h:61
cuda/__memcpy_async/memcpy_async_tx.h:58
cuda/__container/buffer.h:111
cuda/__algorithm/copy.h:73
cuda/__algorithm/fill.h:42
cuda/std/string_view:147

fallback to slower path:

cuda/std/__algorithm/copy.h:102
cuda/std/__algorithm/copy_backward.h:50
cuda/std/__algorithm/move.h:54
cuda/std/__algorithm/move_backward.h:53
cuda/std/__bit/bit_cast.h:55
cub/device/dispatch/kernels/kernel_histogram.cuh:81
cub/detail/uninitialized_copy.cuh:33
thrust/type_traits/is_trivially_relocatable.h:213
cudax/include/cuda/experimental/__kernel/kernel_ref.cuh:54

Require #8347

miscco

I am not too excited about this.

We need to be really careful here because the compiler may also break in some of those cases.

I did some experiments in #3183

bernhardmgruber · 2026-04-01T20:38:11Z

I wanted to start a discussion for some time to add such a trait, but I would really keep this feature internally for now until it is sufficiently baked. I would really like this feature to make thrust::is_trivially_relocatable and THRUST_PROCLAIM_TRIVIALLY_RELOCATABLE obsolete, because those two are often used wrongly, so we should make sure the new trait here can replace it.

davebayer

I am not a fan of this trait. We bend C++ rules to fix poorly designed nvfp types. We would have to basically change every use of is_trivially_copyable and is_trivially_copy_constructible to this trait to make it work consistently in libcu++.

I don't think this is a good idea, we should rather insist of the nvfp types being fixed.

docs/libcudacxx/extended_api/type_traits/is_trivially_copyable_relaxed.rst

davebayer · 2026-04-01T21:12:53Z

docs/libcudacxx/extended_api/type_traits/is_trivially_copyable_relaxed.rst

+Users may specialize ``cuda::is_trivially_copyable_relaxed`` for their own types whose memory representation is safe to copy
+with ``memcpy`` but that the compiler does not consider trivially copyable.


Isn't this just UB?

not for the types that we care about. Said that, the user could provide an object that triggers UB. I can highlight this point in the documentation but we cannot do anything to explicitly prevent it.

fbusato · 2026-04-01T22:26:29Z

We need to be really careful here because the compiler may also break in some of those cases.

I did some experiments in #3183

Integrated and extended the tests that you point out. Everything works

fbusato · 2026-04-01T22:27:40Z

I wanted to start a discussion for some time to add such a trait, but I would really keep this feature internally for now until it is sufficiently baked. I would really like this feature to make thrust::is_trivially_relocatable and THRUST_PROCLAIM_TRIVIALLY_RELOCATABLE obsolete, because those two are often used wrongly, so we should make sure the new trait here can replace it.

the actual issue is that the user can extend the trait to custom types. It needs to be public

fbusato · 2026-04-01T22:29:56Z

I am not a fan of this trait. We bend C++ rules to fix poorly designed nvfp types. We would have to basically change every use of is_trivially_copyable and is_trivially_copy_constructible to this trait to make it work consistently in libcu++.
I don't think this is a good idea, we should rather insist of the nvfp types being fixed.

there are no plan for that. nvfp types are not trivially copyable for optimization purposes.

jrhemstad · 2026-04-01T23:03:39Z

we should rather insist of the nvfp types being fixed.

This isn't going to happen and we shouldn't delude ourselves into thinking it ever will.

libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h

libcudacxx/include/cuda/std/__bit/bit_cast.h

libcudacxx/include/cuda/std/__type_traits/aggregate_members.h

miscco · 2026-04-09T08:08:18Z

libcudacxx/include/cuda/std/__type_traits/aggregate_members.h

Please move this out of the PR.

This is a nontrivial piece of code and the current implementation seems really costly.

I want this in a separate PR so that we can properly review it

Fine, I based the implementation on the code from stdexec. Eric already review it

I am not against merging it, but I want it in a separate PR

created a new PR #8347

libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.aggr.pass.cpp

libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.pass.cpp

Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>

…laxed-type-traits

fbusato · 2026-04-09T18:21:01Z

libcudacxx/include/cuda/std/__bit/bit_cast.h

-#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>,


C++ specification does not impose this constrain https://eel.is/c++draft/bit.cast. is_trivially_default_constructible_v is too strict. cuda::std::complex fails for example.

We only need to check if _To __temp; compiles. The best match is C++20 std::default_initializable

…elaxed-type-traits # Conflicts: # libcudacxx/include/cuda/std/__cccl/builtin.h

miscco · 2026-04-10T09:26:01Z

@fbusato I checked the standard and it turns out that C++23 made complex trivially copyable for standard floating point types, so we should only specialize for __half and __nv_bfloat16

See #8356

…elaxed-type-traits

github-actions · 2026-04-11T02:28:53Z

😬 CI Workflow Results

🟥 Finished in 3h 41m: Pass: 92%/162 | Total: 8d 13h | Max: 3h 41m | Hits: 38%/392424

See results here.

is_trivially_copyable_relaxed

ea3f956

fbusato self-assigned this Apr 1, 2026

fbusato added this to CCCL Apr 1, 2026

fbusato added the libcu++ For all items related to libcu++ label Apr 1, 2026

fbusato requested a review from a team as a code owner April 1, 2026 19:51

github-project-automation bot moved this to Todo in CCCL Apr 1, 2026

fbusato requested a review from wmaxey April 1, 2026 19:51

cccl-authenticator-app bot moved this from Todo to In Review in CCCL Apr 1, 2026

add documentation

9dd60db

fbusato requested a review from a team as a code owner April 1, 2026 19:55

fbusato requested a review from gonidelis April 1, 2026 19:55

fbusato added 2 commits April 1, 2026 13:02

describe custom specialization

179a81b

move to internal function

2c33c2e

miscco requested changes Apr 1, 2026

View reviewed changes

github-project-automation bot moved this from In Review to In Progress in CCCL Apr 1, 2026

davebayer requested changes Apr 1, 2026

View reviewed changes

address padding

fbade34

fbusato added 2 commits April 1, 2026 15:33

fix clang

deb622c

do not handle volatile

a553bb6

fbusato added 2 commits April 1, 2026 16:19

unused-local-typedef

de42a45

fix clang pragma

4e7873d

fbusato moved this from In Progress to In Review in CCCL Apr 1, 2026

This comment has been minimized.

Sign in to view

fbusato added 2 commits April 2, 2026 10:36

simplify conditions

6e5021f

improve documentation

8ff54f4

fbusato added 2 commits April 8, 2026 13:04

add documentation

80b09fa

fix compile warnings/errors

cd776c9

This comment has been minimized.

Sign in to view

miscco requested changes Apr 9, 2026

View reviewed changes

fbusato and others added 4 commits April 9, 2026 09:09

Update libcudacxx/include/cuda/std/__bit/bit_cast.h

d342544

Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>

fix license, inline

d491743

Merge branch 'relaxed-type-traits' of github.com:fbusato/cccl into re…

4da9f32

…laxed-type-traits

add complex support and improve bit_cast

c23d451

fbusato commented Apr 9, 2026

View reviewed changes

move __builtin_structured_binding_size to builtin.h

1065c84

fbusato requested a review from a team as a code owner April 9, 2026 18:31

This comment has been minimized.

Sign in to view

__aggregate_all_of

4450c1d

fbusato mentioned this pull request Apr 9, 2026

__aggregate_all_of Internal utility to recursively apply a predicate to aggregate members #8347

Open

fbusato added 2 commits April 9, 2026 12:44

Merge remote-tracking branch 'origin/aggregate_members_all_of' into r…

fc773e0

…elaxed-type-traits # Conflicts: # libcudacxx/include/cuda/std/__cccl/builtin.h

revert builtin.h

5f1fb0c

This comment has been minimized.

Sign in to view

add missing is_trivially_default_constructible

ce37dab

fbusato requested a review from a team as a code owner April 9, 2026 21:50

fbusato added 2 commits April 9, 2026 16:05

skip MSVC unsupported cases

352b6b3

fix headers

ceb8393

This comment has been minimized.

Sign in to view

fbusato added 2 commits April 10, 2026 11:51

Merge remote-tracking branch 'origin/aggregate_members_all_of' into r…

4b8e70d

…elaxed-type-traits

fix complex and removed user specialization

6813780

This comment has been minimized.

Sign in to view

add half/bfloat guard

6025190

This comment has been minimized.

Sign in to view

protect __half/bfloat16

e432b8d

		Users may specialize ``cuda::is_trivially_copyable_relaxed`` for their own types whose memory representation is safe to copy
		with ``memcpy`` but that the compiler does not consider trivially copyable.

Conversation

fbusato commented Apr 1, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Potentially affected paths by std::is_trivially_copyable

Uh oh!

miscco left a comment

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber commented Apr 1, 2026

Uh oh!

davebayer left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

davebayer Apr 1, 2026

Choose a reason for hiding this comment

Uh oh!

fbusato Apr 1, 2026

Choose a reason for hiding this comment

Uh oh!

fbusato commented Apr 1, 2026

Uh oh!

fbusato commented Apr 1, 2026

Uh oh!

fbusato commented Apr 1, 2026

Uh oh!

jrhemstad commented Apr 1, 2026

Uh oh!

This comment has been minimized.

This comment has been minimized.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

miscco Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

fbusato Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

miscco Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

fbusato Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

fbusato Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

This comment has been minimized.

This comment has been minimized.

This comment has been minimized.

miscco commented Apr 10, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

This comment has been minimized.

This comment has been minimized.

github-actions bot commented Apr 11, 2026

😬 CI Workflow Results

🟥 Finished in 3h 41m: Pass: 92%/162 | Total: 8d 13h | Max: 3h 41m | Hits: 38%/392424

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

fbusato commented Apr 1, 2026 •

edited

Loading

Potentially affected paths by `std::is_trivially_copyable`

miscco commented Apr 10, 2026 •

edited

Loading