gpu sorting algorithms

**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()); 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.