cuda kernel optimization

**CUDA Kernel Optimization** is **the systematic tuning of GPU kernels to maximize throughput, minimize latency, and achieve peak hardware utilization** — where optimizations like memory coalescing (achieving 80-100% memory bandwidth), occupancy tuning (70-100% SM utilization), warp divergence elimination (reducing branch penalties by 50-90%), and instruction-level parallelism (ILP) increase performance by 2-10× over naive implementations through techniques like shared memory tiling that reduces global memory accesses by 80-95%, register optimization that enables 50-100% more active warps, and loop unrolling that improves ILP by 2-4×, making kernel optimization critical for achieving 50-80% of theoretical peak performance (20-40 TFLOPS on A100, 60-80 TFLOPS on H100) where unoptimized kernels typically achieve only 5-20% of peak and systematic optimization following the CUDA performance guidelines can improve performance by 5-20× through memory, compute, and control flow optimizations. **Memory Coalescing:** - **Aligned Access**: threads in warp access consecutive memory addresses; 128-byte aligned; achieves 100% memory bandwidth utilization - **Stride Patterns**: unit stride (consecutive) optimal; stride-2 achieves 50% bandwidth; stride-32 achieves 3% bandwidth; avoid non-unit strides - **Structure of Arrays (SoA)**: prefer SoA over AoS; enables coalesced access; 5-10× memory bandwidth improvement - **Padding**: add padding to avoid bank conflicts; align to 128 bytes; 10-30% performance improvement **Occupancy Optimization:** - **Register Usage**: reduce registers per thread; enables more active warps; 32-64 registers optimal; >128 registers limits occupancy - **Shared Memory**: balance shared memory usage; 48KB per SM on A100; excessive usage reduces occupancy; 16-32KB per block typical - **Block Size**: 128-256 threads per block optimal; too small wastes resources; too large limits occupancy; multiple of 32 (warp size) - **Occupancy Calculator**: use CUDA occupancy calculator; predicts occupancy from resource usage; target 50-100% occupancy **Warp Divergence:** - **Branch Elimination**: remove branches when possible; use arithmetic instead; 2-5× speedup for divergent branches - **Warp-Uniform Branches**: ensure all threads in warp take same path; predicate execution; eliminates divergence penalty - **Thread Coarsening**: assign multiple elements per thread; reduces divergence; 20-50% performance improvement - **Ballot/Shuffle**: use warp-level primitives; avoid explicit synchronization; 2-10× faster than shared memory **Shared Memory Optimization:** - **Tiling**: load data into shared memory; reuse across threads; reduces global memory accesses by 80-95%; 5-20× speedup - **Bank Conflicts**: avoid accessing same bank simultaneously; 32 banks on modern GPUs; stride by 33 to avoid conflicts - **Padding**: add padding to shared memory arrays; prevents bank conflicts; 1-2 elements padding typical - **Synchronization**: minimize __syncthreads(); only when necessary; 10-30% overhead per sync **Register Optimization:** - **Register Pressure**: monitor register usage; nvcc --ptxas-options=-v shows usage; reduce to increase occupancy - **Loop Unrolling**: #pragma unroll; reduces loop overhead; increases ILP; 20-50% speedup; but increases register usage - **Constant Memory**: use __constant__ for read-only data; cached; broadcast to all threads; 2-5× faster than global memory - **Texture Memory**: use texture cache for spatial locality; 2D/3D access patterns; 2-10× speedup for irregular access **Instruction-Level Parallelism:** - **Independent Operations**: reorder instructions; expose ILP; GPU can issue 2-4 instructions per cycle per warp - **Loop Unrolling**: unroll loops by 2-4×; increases ILP; reduces loop overhead; 20-50% speedup - **Multiple Accumulators**: use multiple accumulators in reductions; reduces dependency chains; 30-60% speedup - **Fused Multiply-Add (FMA)**: use FMA instructions; 2× throughput vs separate multiply and add; automatic in most cases **Memory Hierarchy:** - **L1 Cache**: 128KB per SM on A100; automatic caching; prefer shared memory for explicit control - **L2 Cache**: 40MB on A100, 50MB on H100; shared across SMs; benefits from temporal locality - **Global Memory**: 40-80GB HBM2/HBM3; 1.5-3 TB/s bandwidth; minimize accesses; coalesce when accessing - **Unified Memory**: automatic migration; convenient but slower; explicit management preferred for performance **Compute Optimization:** - **Tensor Cores**: use for matrix operations; 312 TFLOPS (FP16) on A100, 989 TFLOPS on H100; 10-20× faster than CUDA cores - **Mixed Precision**: FP16 for compute, FP32 for accumulation; 2× throughput; maintains accuracy; automatic mixed precision (AMP) - **Math Libraries**: use cuBLAS, cuDNN, cuFFT; highly optimized; 2-10× faster than custom kernels - **Warp-Level Primitives**: __shfl, __ballot, __any, __all; faster than shared memory; 2-5× speedup for reductions **Launch Configuration:** - **Grid Size**: enough blocks to saturate GPU; 100-1000 blocks typical; more blocks than SMs for load balancing - **Block Size**: 128-256 threads optimal; multiple of 32; balance occupancy and resource usage - **Dynamic Parallelism**: launch kernels from device; reduces CPU-GPU synchronization; 20-50% overhead; use sparingly - **Streams**: overlap compute and memory transfers; 2-4 streams typical; 20-50% throughput improvement **Profiling Tools:** - **Nsight Compute**: detailed kernel profiling; memory, compute, occupancy metrics; identifies bottlenecks - **Nsight Systems**: timeline view; CPU-GPU interaction; kernel launches, memory transfers; system-level optimization - **nvprof**: command-line profiler; deprecated but still useful; quick performance overview - **Metrics**: achieved occupancy, memory throughput, compute throughput, warp execution efficiency; guide optimization **Common Bottlenecks:** - **Memory Bound**: <50% memory bandwidth; optimize coalescing, use shared memory, reduce accesses - **Compute Bound**: <50% compute throughput; use Tensor Cores, increase ILP, reduce divergence - **Latency Bound**: low occupancy; reduce register usage, increase block size, optimize shared memory - **Instruction Bound**: high instruction overhead; reduce branches, use warp primitives, optimize control flow **Optimization Workflow:** - **Profile**: identify bottleneck; memory, compute, or latency; use Nsight Compute - **Optimize**: apply relevant optimizations; memory coalescing, shared memory, occupancy tuning - **Measure**: verify improvement; compare metrics; iterate if needed - **Iterate**: repeat for next bottleneck; diminishing returns after 3-5 iterations; 2-10× total speedup typical **Advanced Techniques:** - **Cooperative Groups**: flexible thread synchronization; grid-wide sync; warp-level primitives; more expressive than __syncthreads() - **Warp Specialization**: different warps perform different tasks; reduces divergence; 20-40% speedup for heterogeneous workloads - **Persistent Threads**: threads loop over work items; reduces kernel launch overhead; 10-30% speedup for small kernels - **Asynchronous Copy**: async memory copy; overlaps with compute; 20-50% speedup; requires compute capability 8.0+ **Performance Targets:** - **Memory Bandwidth**: 80-100% of peak (1.5-3 TB/s); coalesced access, minimal bank conflicts - **Compute Throughput**: 50-80% of peak (20-40 TFLOPS FP32, 60-80 TFLOPS FP16); use Tensor Cores, high ILP - **Occupancy**: 50-100%; balance register and shared memory usage; 256 threads per block typical - **Warp Efficiency**: >90%; minimize divergence; uniform control flow **Case Studies:** - **Matrix Multiplication**: 80-95% of peak with tiling and Tensor Cores; 10-20 TFLOPS on A100 - **Reduction**: 60-80% of peak with warp primitives and multiple accumulators; 500-1000 GB/s - **Convolution**: 70-90% of peak with cuDNN or custom kernels; 15-30 TFLOPS on A100 - **Sorting**: 40-60% of peak with radix sort; 100-300 GB/s; memory-bound operation **Common Mistakes:** - **Uncoalesced Access**: stride access patterns; 10-100× slowdown; use SoA, align data - **Excessive Synchronization**: too many __syncthreads(); 10-30% overhead each; minimize usage - **Low Occupancy**: too many registers or shared memory; limits active warps; reduce resource usage - **Branch Divergence**: divergent branches within warps; 2-32× slowdown; eliminate or make uniform **Best Practices:** - **Start Simple**: get correct implementation first; then optimize; premature optimization wastes time - **Profile-Guided**: always profile before optimizing; focus on bottlenecks; 80/20 rule applies - **Incremental**: optimize one aspect at a time; measure impact; easier to debug - **Use Libraries**: cuBLAS, cuDNN, Thrust; highly optimized; 2-10× faster than custom code **Performance Portability:** - **Compute Capability**: code for target GPU; A100 (8.0), H100 (9.0); use __CUDA_ARCH__ for conditional compilation - **Tuning Parameters**: block size, tile size, unroll factors; auto-tune for different GPUs; 20-50% performance variation - **Tensor Cores**: available on Volta (7.0) and newer; check capability; fallback to CUDA cores - **Memory Bandwidth**: varies by GPU; A100 (1.5 TB/s), H100 (3 TB/s); adjust algorithms accordingly CUDA Kernel Optimization represents **the art and science of GPU programming** — by applying memory coalescing, occupancy tuning, warp divergence elimination, and shared memory tiling, developers achieve 2-10× performance improvement and 50-80% of theoretical peak performance, making systematic kernel optimization essential for competitive GPU applications where unoptimized kernels achieve only 5-20% of peak and following CUDA best practices can improve performance by 5-20× through memory, compute, and control flow optimizations.');