`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.

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

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

fbusato · 2026-04-30T20:23:42Z

entirely refactored the implementation relying on __aggregate_all_of_v.
replaced cuda::std::is_trivially_copyable with cuda::is_trivially_copyable in all CCCL library paths (but not in tests/benchmarks)

fbusato · 2026-04-30T20:24:29Z

@@ -0,0 +1,323 @@
+//===----------------------------------------------------------------------===//


I'm inclined to remove this file because it is very similar with bit_cast.pass.cpp

Jacobfaib · 2026-05-01T13:23:05Z

+   } // namespace cuda
+
+``cuda::is_trivially_copyable_v`` trait evaluates if a type can be copied by copying its underlying bytes.
+It extends ``cuda::std::is_trivially_copyable`` to also recognize CUDA extended floating-point vector types as trivially copyable.


Should also mention whether it is legal for users to specialize the trait, and give an example of how to do that if so. Given the definition above it implies that users should specialize is_trivially_copyable_v instead of is_trivially_copyable, usually it is the other way around.

specializing is_trivially_copyable_v could be too far. The only reason we are proving it is to overcome CUDA data type limitations. I don't see a good motivation for user to specialize it

@miscco @davebayer let me know if you have a different opinion

I don't see a good motivation for user to specialize it

Even so, I still think the definition snippet should show the "usual" way these traits are defined, i.e.

template <typename T> struct my_trait : bool_constant<...> { }; template <typename T> inline constexper bool my_trait_v = my_trait<T>::value;

Even if that is not technically true, we don't want to give the impression that specializing the _v form is what actually matters.

Jacobfaib · 2026-05-01T13:23:40Z

+``cuda::is_trivially_copyable_v`` trait evaluates if a type can be copied by copying its underlying bytes.
+It extends ``cuda::std::is_trivially_copyable`` to also recognize CUDA extended floating-point vector types as trivially copyable.
+
+``cuda::is_trivially_copyable_v<T>`` relies on ``cuda::std::is_trivially_copyable`` but adds support for CUDA-specific types.


If it is legal to specialize, should elaborate whether the user should specialize cuda::std::is_trivially_copyable or cuda::is_trivially_copyable

Jacobfaib · 2026-05-01T13:50:43Z

Important: can you add tests that the trait works even when it is user-specialized?

pciolkosz · 2026-05-01T22:57:59Z

+``cuda::is_trivially_copyable_v`` trait evaluates if a type can be copied by copying its underlying bytes.
+It extends ``cuda::std::is_trivially_copyable`` to also recognize CUDA extended floating-point vector types as trivially copyable.
+
+``cuda::is_trivially_copyable_v<T>`` relies on ``cuda::std::is_trivially_copyable`` but adds support for CUDA-specific types.


Note: This sentence and the previous one seem a bit redundant

github-actions · 2026-05-02T01:48:38Z

😬 CI Workflow Results

🟥 Finished in 3h 26m: Pass: 97%/437 | Total: 19d 02h | Max: 3h 25m | Hits: 44%/599270

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

protect __half/bfloat16

e432b8d

This comment has been minimized.

Sign in to view

fbusato moved this from In Progress to Blocked in CCCL Apr 14, 2026

fbusato moved this from Blocked to In Progress in CCCL Apr 28, 2026

fbusato moved this from In Progress to Todo in CCCL Apr 28, 2026

Merge branch 'main' into relaxed-type-traits

6f49e7b

This comment has been minimized.

Sign in to view

fbusato and others added 6 commits April 30, 2026 11:47

align with is_bitwise_comparable

55b641f

Merge branch 'main' into relaxed-type-traits

d328eb1

update implementation and unit tests

26f9e6c

apply cuda::is_trivially_copyable to all files

3c62e3f

split bit_cast unit test

4bc8737

formatting

38d3419

fbusato requested review from a team as code owners April 30, 2026 20:10

update documentation

9e76b3f

fbusato requested a review from miscco April 30, 2026 20:21

fbusato commented Apr 30, 2026

View reviewed changes

This comment has been minimized.

Sign in to view

Jacobfaib reviewed May 1, 2026

View reviewed changes

fbusato and others added 7 commits May 1, 2026 13:18

fix CI failures

c5d1c4a

add TEST_FUNC

a44c391

update doc

7addbc1

pre-commit

7fefe97

exclude default_initializable for GCC7

b40820c

Merge branch 'main' into relaxed-type-traits

eed543a

fix doc

55402fe

pciolkosz reviewed May 2, 2026

View reviewed changes

		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.

		@@ -0,0 +1,323 @@
		//===----------------------------------------------------------------------===//

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!

Choose a reason for hiding this comment

Uh oh!

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.

This comment has been minimized.

fbusato commented Apr 30, 2026

Uh oh!

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

Uh oh!

Choose a reason for hiding this comment

Uh oh!

This comment has been minimized.

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented May 2, 2026

😬 CI Workflow Results

🟥 Finished in 3h 26m: Pass: 97%/437 | Total: 19d 02h | Max: 3h 25m | Hits: 44%/599270

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

8 participants

fbusato commented Apr 1, 2026 •

edited

Loading

Potentially affected paths by `std::is_trivially_copyable`

fbusato Apr 30, 2026 •

edited

Loading