NVIDIA / NVIDIA/cccl

Refine fallback kernel for segmented sort

Open
#895 1 comment 0 reactions 0 assignees View on GitHub
cub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.