NVIDIA / NVIDIA/cccl

[FEA]: Let block-algorithms take const arrays for input parameters

Open
#1,296 0 comments 1 reaction 0 assignees View on GitHub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.