Refine fallback kernel for segmented sort
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
Currently, `cub::DeviceSegmentedSort` has a fallback kernel, that [apply](https://github.com/NVIDIA/cub/blob/0430cc0bfcb7c2496b42da754c215c9b5df8856b/cub/device/dispatch/dispatch_segmented_sort.cuh#L169) different algorithms for different segment sizes. In particular, medium-size segments are sorted by merge sort. If segment doesn't fit into registers, it's sorted by in-shared-memory block radix sort. Otherwise, expensive in-global-memory radix sort is applied.
The suggestion was made, that we have to add small-segment branch and potentially experiment with refining the granularity at which we assign threads to segments at fallback kernel.
Contributor guide
Research direction
Read cub/device/dispatch/dispatch_segmented_sort.cuh around the DeviceSegmentedSort fallback kernel at line 169, then trace how segment sizes select the existing algorithms. Determine how to add the suggested small-segment branch and evaluate finer thread-to-segment granularity; done means the fallback strategy has been refined based on those experiments.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- performance
- Issue type
- Refactor
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100