[FEA]: Provide a way to create an uninitialized array

[miscco]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

CUB

Is your feature request related to a problem? Please describe.

For vectorized code paths we often use something like InputT input[items_per_thread];

Because this is a plain C-array, the compiler will try to value-initialize the elements. However, this is almost never what we want because we want to fill it with data usually through a vectorized load.

Describe the solution you'd like

We should provide a type that holds the array in a union, so that we can create it without spending time on initializing it.

Describe alternatives you've considered

Alternatively we could use alignas(InputT) cuda::std::byte input[items_per_thread * sizeof(InputT)]; and then cast, but I am not sure what is cleaner

Additional context

No response

[fbusato]

Could you provide an example of when the compiler cannot understand that the array has not been overwritten?

[miscco]

It starts with stuff like this #7336

[fbusato]

I misunderstood the issue. This is a C++ initialization problem, not a missed nvcc optimization.

[bernhardmgruber]

We do have cub::Uninitialized<T> but it serves a slightly different purpose:

cccl/cub/cub/util_type.cuh

Line 645 in 6e592be

struct Uninitialized

And then we have thrust::cuda_cub::core::detail::uninitialized_array<T, B, Alignment = alignof(T)> which is probably exactly what you are looking for:

cccl/thrust/thrust/system/cuda/detail/core/util.h

Line 514 in b7ebe1e

struct uninitialized_array