[FEA]: All relevant CUB block wide collective primitives should optionally take a valid_items parameter.
- 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.
Often times when I use CUB histogram I need to compensate for valid items due to thread over-provisioning. This field exists on CUB blockwide merge sort, so in theory I don't see any technical reason why it shouldn't be usable in primitives like `cub::BlockHistogram(...)`. With out this I have to manually calculate the amount of extra items per thread, then subtract that value from the histogram from a single thread over the entire block, and put a `__syncthreads()` to compensate. If there is a technical reason a `valid_items` overload cannot be used, then at the very least how to work around this use case should be common enough to be documented explicitly (in the cub examples repository for example). I believe the only two relevant CUB block primitives that don't implement `valid_items` behavior are `cub::BlockHistogram` and `cub::BlockScan`.
### Describe the solution you'd like
The following overloads should be added to `cub::BlockHistogram`
```
template
inline void Composite(
T (&items)[ItemsPerThread],
CounterT histogram[Bins]
int valid_items
)
inline void Histogram(
T (&items)[ItemsPerThread],
CounterT histogram[Bins]
int valid_items
)
```
And similar overloads for `cub::BlockScan`
### Describe alternatives you've considered
Manually accounting for the histogram offsets induces additional overhead that i don't believe is necessary.
### Additional context
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.