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

[pauleonix]

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.

#include <cub/cub.cuh>

__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<int, 128> 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