[FEA]: Let block-algorithms take const arrays for input parameters
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
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.
```cuda
#include
__global__ void foo() {
const int tid = blockDim.x * blockIdx.x + threadIdx.x;
const int gstride = blockDim.x * gridDim.x;
const int thread_data[] = {tid, tid + gstride};
typedef cub::BlockReduce BlockReduce;
__shared__ typename BlockReduce::TempStorage temp_storage;
const int aggregate = BlockReduce(temp_storage).Sum(thread_data);
}
```
does not compile because the interfaces by CUB take non-const arrays as input parameters. This prevents users from writing const-correct kernel functions using CUB block-scope algorithms. `cub::BlockReduce::Sum` is only one example. The same issue appears for basically all interfaces taking multiple input values as a C-style array including `cub::StoreDirectBlocked` and similar direct store functions.
### Describe the solution you'd like
As confirmed by @gevtushenko in #1288, CUB's block-scope algorithms (in particular the overloads taking C arrays residing in registers or local memory) should take const-qualified arrays for input parameters.
### Describe alternatives you've considered
Not writing const-correct kernel-functions can lead to bugs and makes code less readable.
### Additional context
_No response_
Contributor guide
Research direction
Start by inspecting the CUB block-scope algorithm interfaces, especially cub::BlockReduce::Sum and the direct-store functions such as cub::StoreDirectBlocked. Check how their C-style array input parameters are declared and identify the relevant tests or compile examples. Done means the shown const array kernel and equivalent block-algorithm inputs compile with const-qualified arrays.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- backend-api-design
- Issue type
- Feature
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 45/100