[FEA]: Add LRB segmented-sort kernels 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 needs kernels for sub-warp/warp, small-block, and large-block work-plan tiers. These kernels have similar bodies and use the same runtime-width group merge-sort primitive, but require different execution and resource configurations.
### Describe the solution you'd like
Track the three tier-specific kernels:
- [ ] #10947: sub-warp/warp kernel.
- [ ] #10948: small-block kernel.
- [ ] #10949: large-block kernel for segments that fit within one CTA-scale tile.
Each kernel should consume its work-plan tier and invoke the runtime-width group merge-sort primitive tracked in #10944.
This issue can be closed once optimized segmented sort algorithms can be tested and benchmarked but not reachable by user.
cub::detail::lrb_based_segmented_sort
### Describe alternatives you've considered
A single kernel for all tiers was considered, but the tiers require different execution and resource configurations.
### Additional context
Contributor guide
Research direction
Start at cub::detail::lrb_based_segmented_sort and review the runtime-width group merge-sort primitive tracked in #10944, along with the tier-specific work in #10947, #10948, and #10949. Done means the sub-warp/warp, small-block, and large-block kernels consume their work-plan tiers and the optimized algorithms can be tested and benchmarked without being reachable by users.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc
- Issue type
- Feature
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 45/100