NVIDIA / NVIDIA/cccl

[FEA]: All relevant CUB block wide collective primitives should optionally take a valid_items parameter.

Open
#8,202 3 comments 0 reactions 1 assignee Claimed by @miscco View on GitHub
needs triage
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

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.