[FEA]: Add a large-block segmented-sort kernel using the runtime-width group merge primitive
- 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
Assessment
This issue has not been assessed yet.