NVIDIA / NVIDIA/cccl

[FEA]: Add a large-block segmented-sort kernel using the runtime-width group merge primitive

Open
#10,949 1 comment 0 reactions 1 assignee Claimed by @kaatish 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.

Fine-grained segmented-sort planning can assign 512 or 1024 cooperating threads to a segment that fits within a CTA-scale tile. A kernel is needed to consume these large-block work-plan configurations.

### Describe the solution you'd like

Add a large-block segmented-sort kernel that:

- Consumes work-plan configurations using 512 or 1024 threads per segment.
- Uses the runtime-width group merge-sort primitive tracked in #10944.
- Supports the shared-memory requirements and partially filled tiles for these configurations.
- Provides correctness coverage for the supported widths and sort modes.

Segments exceeding the supported CTA-scale capacity remain the responsibility of the fallback or separate large-segment load-balancing work.

This issue can be closed once the kernel is implemented in cub::detail and tested as a consumer of the large-block work-plan tier.

### Describe alternatives you've considered

Separate sorting implementations for each width are unnecessary because they share the runtime-width primitive.

### Additional context

Parent tracking issue: #10946.

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.