NVIDIA / NVIDIA/cccl

[FEA]: Make CUB block algorithms usable with cuda::std::array

Open
#1,877 16 comments 1 reaction 1 assignee Claimed by @pauleonix View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

### Is this a duplicate?

- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Area

CUB

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

#1764 deprecated `cub::ArrayWrapper` in favor of `cuda::std::array` without providing a replacement for users using it with CUB's block algorithms that take C-style arrays. `cub::ArrayWrapper` provided access to it's C-style `array` member, but `cuda::std::array` doesn't (because `std::array` doesn't either). Inside CUB this is solved by accessing `cuda::std::arrays`'s `.__elems_` (an implementation detail) which is not a good solution for users.

### Describe the solution you'd like

~~Replace the C-style arrays in CUBs interfaces with `cuda::std::span`s.~~ Use fbusato's "minimal 'concepts' for array-like types" as proposed in #2286.

### Describe alternatives you've considered

`cuda::std::span` seems like the right candidate for the interfaces but with `T` being a template type in the CUB algorithm [this would still need a duplication of the interfaces](https://stackoverflow.com/a/70118572/10107454) to stay backward compatible.

As C-style arrays decay to pointers when returned from functions, a simple function like `cub::to_array(cuda::std::array)` seems impossible. It could be done with a macro that accesses `.__elems_` but macros have their own problems.

Adding an API to `cuda::std::array` that is not available for `std::array` is probably against libcu++'s principles.

Adding a whole new `cuda::array` for access to the C-style array member seems over the top ~~although it might still be easier to maintain than adding overloads taking `cuda::std::array` parameters to all CUB block algorithms.~~ (Edit: This can be solved using `cuda::std::span` for the interface as pointed out by miscco below)

I am currently using my own array wrapper to avoid plain C-style arrays in my code when interfacing with CUB's block algorithms.

### Additional context

_No response_

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.