NVIDIA / NVIDIA/cccl

[FEA]: Load-balanced segmented reduce

Open
#6,171 1 comment 1 reaction 0 assignees View on GitHub
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.

I would like better performance from `cub::DeviceSegmentedReduce` for "small" segment sizes O(1)/O(10).

The current implementation uses a simple mapping of 1 CTA per segment which is inefficient when segment sizes are small like this.

### Describe the solution you'd like

I would like `cub::DeviceSegmentedReduce` to take advantage of the same "load balancing" approach the @gevtushenko developed for `cub::DeviceSegmentedSort` that involves binning segments based on size and dispatching to different levels of parallelism depending on how big the segments are.

@gevtushenko already had a PoC implementation in the old repo that showed dramatic performance improvements https://github.com/NVIDIA/cub/pull/578

Image

This issue can be closed by picking up where https://github.com/NVIDIA/cub/pull/578 left off and refactor `cub::DeviceSegmentedReduce` to take advantage of the load balanced approach with improved performance for small segment sizes.

### Describe alternatives you've considered

_No response_

### 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.