[FEA]: Improve cub::DeviceHistogram performance
- 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.
General tracking issue for various optimizations that impact or target the high level `cub::DeviceHistogram` on multiple ends. Optimizations could be in any level (block, warp etc.)
- [x] [Use atomicAdd_block in device histogram](#4973)
- [ ] Use warp atomic aggregates
- [ ] Experiment with saturation small-sized counters
- [x] Evaluate `__match_sync_any` to replace warp-level atomic operations (slower)
### Describe the solution you'd like
Overall performance improvement of `cub::*Histogram`
### Describe alternatives you've considered
_No response_
### Additional context
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.