cuda::is_trivially_copyable_relaxed

[fbusato]

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

[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]

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]

Isn't this just UB?

[fbusato]

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]

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]

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]

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]

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.

[github-actions]

😬 CI Workflow Results

🟥 Finished in 1h 18m: Pass: 90%/99 | Total: 1d 03h | Max: 1h 12m | Hits: 99%/231754

See results here.