[FEA]: Investigate WarpRadixSorting
- 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
Assessment
This issue has not been assessed yet.