GPU Sorting Algorithms are the parallel implementations of sorting that leverage thousands of GPU threads to achieve 100-300 GB/s throughput — where radix sort (optimal for integers and fixed-point) achieves 200-300 GB/s by processing multiple bits per pass and exploiting warp-level primitives, merge sort (optimal for general comparisons) achieves 100-200 GB/s through hierarchical merging, and bitonic sort (optimal for power-of-2 sizes) achieves 150-250 GB/s with fixed communication patterns, making GPU sorting 10-50× faster than CPU sorting (5-20 GB/s) and essential for applications like database operations, graph algorithms, and data preprocessing where sorting is bottleneck (20-60% of runtime) and proper algorithm selection based on data characteristics (integer vs float, key-only vs key-value, size) determines whether applications achieve 40% or 90% of theoretical peak bandwidth.
Radix Sort:
- Algorithm: sorts by processing k bits per pass; typically k=4-8 bits; requires ceil(32/k) passes for 32-bit integers; stable sort
- Performance: 200-300 GB/s on A100; 60-80% of peak memory bandwidth; optimal for integers, fixed-point; 10-50× faster than CPU
- Implementation: histogram per block → prefix sum → scatter; uses shared memory for local histogram; warp primitives for reduction
- Use Cases: integer keys, fixed-point values; uniform distribution; large arrays (>1M elements); 80-95% of peak bandwidth
Merge Sort:
- Algorithm: hierarchical merging; bottom-up or top-down; log2(N) passes; stable sort; comparison-based
- Performance: 100-200 GB/s on A100; 40-60% of peak bandwidth; optimal for general comparisons, small arrays
- Implementation: warp-level merge → block-level merge → global merge; uses shared memory for local merging
- Use Cases: general comparisons, custom comparators; small-medium arrays (10K-1M elements); stable sort required
Bitonic Sort:
- Algorithm: comparison network; fixed communication pattern; log2(N) × (log2(N)+1) / 2 comparisons; not stable
- Performance: 150-250 GB/s on A100; 50-70% of peak bandwidth; optimal for power-of-2 sizes; predictable performance
- Implementation: warp-level bitonic → block-level bitonic → global bitonic; uses shuffle for warp-level, shared memory for block-level
- Use Cases: power-of-2 sizes, predictable latency; small-medium arrays; GPU-friendly communication pattern
Thrust Sort:
- API: thrust::sort(d_vec.begin(), d_vec.end()); automatic algorithm selection; radix sort for integers, merge sort for general
- Performance: 100-300 GB/s; 60-80% of hand-tuned; 1 line of code vs 100-200 for custom implementation
- Customization: thrust::sort(d_vec.begin(), d_vec.end(), thrust::greater<int>()); custom comparators supported
- Use Cases: rapid development, general-purpose sorting; acceptable 10-30% performance gap vs hand-tuned
CUB Sort:
- API: cub::DeviceRadixSort::SortKeys(d_temp_storage, temp_storage_bytes, d_keys, d_sorted, N); explicit control
- Performance: 200-300 GB/s; 70-90% of peak bandwidth; 10-30% faster than Thrust; lower-level API
- Features: key-only or key-value pairs; ascending or descending; segmented sort; double-buffer for in-place
- Use Cases: performance-critical sorting; fine-grained control; production systems
Key-Value Sorting:
- Radix Sort: cub::DeviceRadixSort::SortPairs(); sorts keys, reorders values; 150-250 GB/s; 2× slower than key-only
- Merge Sort: stable sort preserves value order; 80-150 GB/s; 40-60% of peak bandwidth
- Performance: 2-3× slower than key-only; memory bandwidth limited; 2× data movement
- Use Cases: sorting with associated data; database operations; graph algorithms
Segmented Sort:
- Concept: sort multiple independent segments; each segment sorted separately; useful for batched operations
- Implementation: cub::DeviceSegmentedRadixSort::SortKeys(); segment offsets specify boundaries
- Performance: 150-250 GB/s; 50-70% of peak; depends on segment sizes; small segments have overhead
- Use Cases: batched sorting, per-group sorting; graph algorithms; database operations
Optimization Techniques:
- Warp-Level Primitives: use __shfl for warp-level sorting; 2-5× faster than shared memory; 500-1000 GB/s for small arrays
- Shared Memory: use for block-level sorting; 100× faster than global memory; 164KB per SM on A100
- Coalesced Access: ensure coalesced memory access; 128-byte aligned; achieves 100% bandwidth; stride-1 access optimal
- Occupancy: balance shared memory usage and occupancy; 50-100% occupancy typical; 256 threads per block optimal
Radix Sort Optimization:
- Bits Per Pass: 4-8 bits typical; more bits = fewer passes but larger histogram; 4 bits optimal for most cases
- Histogram: per-block histogram in shared memory; warp primitives for reduction; 300-600 GB/s
- Prefix Sum: exclusive scan of histogram; 400-800 GB/s; CUB provides optimized implementation
- Scatter: write sorted elements to output; coalesced writes; 200-300 GB/s; double-buffering eliminates copies
Merge Sort Optimization:
- Warp-Level Merge: use __shfl for merging within warp; 2-5× faster than shared memory; 500-1000 GB/s
- Block-Level Merge: shared memory for merging within block; 100-200 GB/s; minimize global memory accesses
- Global Merge: hierarchical merging; multiple passes; 80-150 GB/s; memory-bound
- Path Decomposition: parallel merge path algorithm; better load balancing; 20-40% speedup
Bitonic Sort Optimization:
- Warp-Level: use __shfl_xor for butterfly exchanges; 2-5× faster than shared memory; 500-1000 GB/s
- Block-Level: shared memory for larger bitonic networks; 150-250 GB/s; minimize bank conflicts
- Padding: add padding to avoid bank conflicts; 1-2 elements typical; 10-30% improvement
- Unrolling: unroll inner loops; reduces overhead; 10-20% speedup; compiler often does automatically
Performance Comparison:
- Radix Sort: 200-300 GB/s; best for integers; 60-80% of peak; 10-50× faster than CPU
- Merge Sort: 100-200 GB/s; best for general comparisons; 40-60% of peak; 5-20× faster than CPU
- Bitonic Sort: 150-250 GB/s; best for power-of-2 sizes; 50-70% of peak; 10-30× faster than CPU
- Thrust Sort: 100-300 GB/s; automatic selection; 60-80% of hand-tuned; easiest to use
Size Considerations:
- Small Arrays (<10K): bitonic sort or warp-level sort; 150-250 GB/s; low overhead; predictable latency
- Medium Arrays (10K-1M): merge sort or radix sort; 150-250 GB/s; good balance; algorithm depends on data type
- Large Arrays (>1M): radix sort for integers, merge sort for general; 200-300 GB/s; amortizes overhead; optimal performance
- Very Large (>100M): out-of-core sorting; multiple passes; 100-200 GB/s; memory-limited
Stability:
- Stable: radix sort, merge sort; preserves relative order of equal elements; required for some applications
- Unstable: bitonic sort, quicksort; may reorder equal elements; faster but less predictable
- Use Cases: stable required for multi-key sorting, database operations; unstable acceptable for unique keys
Custom Comparators:
- Merge Sort: supports custom comparators; thrust::sort(d_vec.begin(), d_vec.end(), my_comparator()); flexible
- Radix Sort: limited to integer keys; can use custom bit extraction; less flexible but faster
- Performance: custom comparators 10-30% slower; function call overhead; inline when possible
Profiling and Tuning:
- Nsight Compute: shows memory bandwidth, occupancy, warp efficiency; identifies bottlenecks
- Metrics: achieved bandwidth / peak bandwidth; target 60-80% for sorting; memory-bound operation
- Bottlenecks: uncoalesced access, bank conflicts, low occupancy; optimize access patterns
- Tuning: adjust block size, bits per pass, shared memory usage; profile to find optimal
Best Practices:
- Use Libraries: Thrust or CUB for most cases; 60-90% of hand-tuned; 10-100× less code
- Algorithm Selection: radix sort for integers, merge sort for general; bitonic for power-of-2; profile to verify
- Batch Operations: sort multiple arrays together; amortizes overhead; 20-40% improvement
- Profile: measure actual bandwidth; compare with peak; optimize only if bottleneck
- Pre-Allocate: allocate temporary storage once; reuse across sorts; eliminates allocation overhead
Performance Targets:
- Radix Sort: 200-300 GB/s; 60-80% of peak (1.5-3 TB/s); optimal for integers
- Merge Sort: 100-200 GB/s; 40-60% of peak; optimal for general comparisons
- Bitonic Sort: 150-250 GB/s; 50-70% of peak; optimal for power-of-2 sizes
- Key-Value: 150-250 GB/s; 50-70% of peak; 2× slower than key-only
Real-World Applications:
- Database Operations: sorting query results; 200-300 GB/s with radix sort; 10-50× faster than CPU
- Graph Algorithms: sorting edges by source/destination; 150-250 GB/s; 20-40% of graph processing time
- Data Preprocessing: sorting features for ML; 200-300 GB/s; 10-30% of preprocessing time
- Rendering: sorting primitives by depth; 150-250 GB/s; 5-20% of rendering time
GPU Sorting Algorithms represent the essential building block for data-intensive applications — by leveraging thousands of parallel threads and optimized memory access patterns, GPU sorting achieves 100-300 GB/s throughput (10-50× faster than CPU) through algorithms like radix sort for integers (200-300 GB/s), merge sort for general comparisons (100-200 GB/s), and bitonic sort for power-of-2 sizes (150-250 GB/s), making GPU sorting critical for applications where sorting is bottleneck and proper algorithm selection based on data characteristics determines whether applications achieve 40% or 90% of theoretical peak bandwidth.