Skip to content

GPU: Faster QuickSort implementation for large data, and some cleanup of option names and help / debug messages #15933

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
davidrohr merged 2 commits into from
Oct 11, 2026
+64 −51
Merged

GPU: Faster QuickSort implementation for large data, and some cleanup of option names and help / debug messages #15933

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
05f62ce
GPU: Fix some names and descriptions in GPU Settings
davidrohr Oct 11, 2026
5686fda
GPU Common Algorithm: Improve QuickSort implementation
davidrohr Oct 11, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
38 changes: 22 additions & 16 deletions GPU/Common/GPUCommonAlgorithm.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,37 +41,34 @@ class GPUCommonAlgorithm
GPUd() static void sortInBlock(T* begin, T* end, const S& comp GPUCA_THREAD_INFO_DECL);
template <class T, class S>
GPUd() static void sortDeviceDynamic(T* begin, T* end, const S& comp);
#if __cplusplus >= 202002L // sortOnDevice takes an auto parameter
#if __cplusplus >= 202002L // FIXME: auto parameter not supported in C++17
template <class T, class S>
GPUh() static void sortOnDevice(auto* rec, int32_t stream, T* begin, size_t N, const S& comp);
#endif
template <class T>
GPUd() static void swap(T& a, T& b);

private:
// Quicksort implementation
template <typename I>
GPUd() static void QuickSort(I f, I l) GPUnoexcept();
GPUd() static void QuickSort(I f, I l) GPUnoexcept(); // Quicksort implementation

// Quicksort implementation
template <typename I, typename Cmp>
GPUd() static void QuickSort(I f, I l, Cmp cmp) GPUnoexcept();
GPUd() static void QuickSort(I f, I l, Cmp cmp) GPUnoexcept(); // Quicksort implementation

template <typename IndexType, typename I, typename Cmp>
GPUd() static void QuickSortImpl(I f, I l, Cmp cmp) GPUnoexcept(); // Quicksort implementation with index type IndexType, which must be able to hold the number of elements

// Insertionsort implementation
template <typename I, typename Cmp>
GPUd() static void InsertionSort(I f, I l, Cmp cmp) GPUnoexcept();
GPUd() static void InsertionSort(I f, I l, Cmp cmp) GPUnoexcept(); // Insertionsort implementation

// Helper for Quicksort implementation
template <typename I, typename Cmp>
GPUd() static I MedianOf3Select(I f, I l, Cmp cmp) GPUnoexcept();
GPUd() static I MedianOf3Select(I f, I l, Cmp cmp) GPUnoexcept(); // Helper for Quicksort implementation

// Helper for Quicksort implementation
template <typename I, typename T, typename Cmp>
GPUd() static I UnguardedPartition(I f, I l, T piv, Cmp cmp) GPUnoexcept();
GPUd() static I UnguardedPartition(I f, I l, T piv, Cmp cmp) GPUnoexcept(); // Helper for Quicksort implementation

// Helper
template <typename I>
GPUd() static void IterSwap(I a, I b) GPUnoexcept();
GPUd() static void IterSwap(I a, I b) GPUnoexcept(); // Helper
};

#ifndef GPUCA_ALGORITHM_STD
Expand Down Expand Up @@ -147,19 +144,28 @@ GPUdi() I GPUCommonAlgorithm::UnguardedPartition(I f, I l, T piv, Cmp cmp) GPUno

template <typename I, typename Cmp>
GPUdi() void GPUCommonAlgorithm::QuickSort(I f, I l, Cmp cmp) GPUnoexcept()
{
if (l - f <= 0xFFFF) {
QuickSortImpl<uint16_t>(f, l, cmp);
} else {
QuickSortImpl<uint32_t>(f, l, cmp);
}
}

template <typename IndexType, typename I, typename Cmp>
GPUdi() void GPUCommonAlgorithm::QuickSortImpl(I f, I l, Cmp cmp) GPUnoexcept()
{
if (f == l) {
return;
}
using IndexType = uint16_t;

struct pair {
IndexType first;
IndexType second;
};

struct Stack {
pair data[11];
pair data[sizeof(IndexType) * 8 - 4]; // The smaller partition is processed first, so the depth is bounded by log2(n / cutoff) + 1, cutoff = 50 > 2^5
uint8_t n{0};

GPUd() void emplace(IndexType x, IndexType y)
Expand Down Expand Up @@ -282,7 +288,7 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp GPU
{
#ifndef GPUCA_GPUCODE
GPUCommonAlgorithm::sort(begin, end, comp);
#elif defined(GPUCA_DETERMINISTIC_MODE) // Not using GPUCA_DETERMINISTIC_CODE, which is enforced in TPC compression
#elif defined(GPUCA_DETERMINISTIC_MODE) // Cannot use GPUCA_DETERMINISTIC_CODE macro here, which is enforced in TPC compression
if (get_local_id(0) == 0) {
GPUCommonAlgorithm::sort(begin, end, comp);
}
Expand Down
Loading
Loading