NVIDIA / NVIDIA/cccl

[FEA]: Investigate WarpRadixSorting

Open
#6,793 1 comment 1 reaction 1 assignee Claimed by @gonidelis View on GitHub
needs triage
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.

Merge-sort and Bitonic-sort generally require power-of-two input sizes, incurring significant overhead for non-conforming inputs. In the context of segmented sorts, this results in substantial wasted work when handling many small segments (e.g., sub-warp sizes). Additionally, because they are comparison-based, they lack the efficiency of radix sorts when handling keys smaller than 32 bits.

### Describe the solution you'd like

Investigate the performance of a WarpRadixSort primitive for its use in Segmented Sorts, in conjunction with next-fit-bin packing:

```cuda
template
__device__ __forceinline__ void MultiSplit32AsmGe(uint32_t& geMask, const uint32_t key) {
#pragma unroll
for (uint32_t bit = 0; bit < BITS_TO_SORT; ++bit) {
uint32_t current_bit = 1 << bit;
asm("{\n"
" .reg .pred p;\n"
" and.b32 %2, %1, %2;\n"
" setp.eq.u32 p, %2, 0;\n"
" vote.ballot.sync.b32 %2, p, 0xffffffff;\n"
" @p and.b32 %0, %0, %2;\n"
" @!p or.b32 %0, %0, %2;\n"
"}\n"
: "+r"(geMask)
: "r"(key), "r"(current_bit));
}
}

template
__device__ __forceinline__ uint32_t Sort32(const uint32_t key) {
uint32_t geMask = getLaneMaskLt();
MultiSplit32AsmGe(geMask, key);

return __popc(geMask);
}

```

### Describe alternatives you've considered

_No response_

### Additional context

_No response_

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.