The goal of this project was to efficiently parallelize Radix Sort on GPUs using CUDA. The design of the algorithm follows the paper Designing Efficient Sorting Algorithms for Manycore GPUs from the authors Satish, Harris a Garland. I made this as my semestral project for the subject GPU Architectures and Programming.
The full detailed report, written in Czech, can be found in docs.
But briefly, each pass of the algorithm consists of 3 kernels (+ Host <-> Device mem transfers etc.).
blockLocalSortAndComputeCountsKernel- Locally (stably) sort each block bit-by-bit in shared memory using a split operation.
- Compute local offsets/starts of each digit.
- Compute the counts of each digit - histogram.
prefixSumBucketCounts- Performs a parallel exclusive prefix sum over the histograms.globalScatterKernel- Scatters the numbers to the global output array, using the result of the PPS + local offsets.
I performed several benchmarks to try to optimize the algorithm futher and to compare it with its sequential version and the highly optimized cub::DeviceRadixSort.
The benchmarks were run on an Apple M4 Pro CPU (14 cores, 48 GB RAM) and an NVIDIA GeForce RTX 4070 Ti GPU.
The tested distributions were uniform (random values), nearlySorted (almost sorted values with small noise), and heavyHitter (many repeated values from a small set).
This graph measures CUDA configuration performance. The x-axis shows ITEMS_PER_THREAD / THREADS_PER_BLOCK, the y-axis shows RADIX_BITS, and each value is the relative runtime compared to the best configuration for the same dataset.
This graph measures the speedup of the best CUDA configuration over the M4 CPU implementation. The x-axis shows the number of sorted elements, and the y-axis shows speedup in x.
This graph measures how much slower the best CUDA configuration was compared to cub::DeviceRadixSort. The x-axis shows the number of sorted elements, and the y-axis shows the runtime ratio in x. As you can see, I was not faster :(