gpu profiling debugging

**GPU Profiling and Debugging** is **the systematic analysis of GPU application performance and correctness using specialized tools that provide detailed metrics, timeline visualization, and error detection** — where NVIDIA Nsight Compute delivers kernel-level analysis with 1000+ metrics covering memory bandwidth (achieved vs peak 1.5-3 TB/s), compute throughput (achieved vs peak 20-80 TFLOPS), occupancy (50-100%), and warp efficiency (target >90%), while Nsight Systems provides system-wide timeline showing CPU-GPU interaction, kernel launches, memory transfers, and API calls, enabling developers to identify bottlenecks (memory-bound, compute-bound, latency-bound), optimize resource utilization, and achieve 2-10× performance improvement through data-driven optimization, making profiling essential for GPU development where intuition often misleads and measurement is the only path to understanding actual performance characteristics. **Nsight Compute (Kernel Profiling):** - **Purpose**: detailed single-kernel analysis; 1000+ metrics; memory, compute, occupancy, warp efficiency; identifies kernel bottlenecks - **Launch**: ncu ./app or ncu --set full ./app; GUI: ncu-ui; command-line or graphical interface - **Metrics**: Memory Throughput (GB/s), Compute Throughput (TFLOPS), SM Efficiency (%), Occupancy (%), Warp Execution Efficiency (%), Branch Efficiency (%) - **Sections**: Memory Workload Analysis, Compute Workload Analysis, Launch Statistics, Occupancy, Scheduler Statistics, Warp State Statistics **Nsight Systems (System Profiling):** - **Purpose**: system-wide timeline; CPU-GPU interaction; kernel launches, memory transfers, API calls; identifies system-level bottlenecks - **Launch**: nsys profile ./app or nsys profile --trace=cuda,nvtx ./app; generates .qdrep file; open in nsys-ui - **Timeline View**: visualizes all GPU activity; shows overlaps, gaps, synchronization points; identifies idle time - **Use Cases**: multi-GPU profiling, stream concurrency, CPU-GPU overlap, kernel launch overhead, memory transfer analysis **Memory Profiling:** - **Memory Throughput**: achieved bandwidth / peak bandwidth; target 80-100% for memory-bound kernels; A100: 1.5-2 TB/s, H100: 2-3 TB/s - **Memory Replay**: indicates uncoalesced access; replay >1.5 means poor coalescing; restructure data layout - **L1/L2 Hit Rate**: cache effectiveness; high hit rate (>80%) good for reused data; low hit rate indicates streaming access - **Global Load/Store Efficiency**: percentage of useful bytes loaded; low efficiency (<50%) indicates wasted bandwidth; improve coalescing - **Bank Conflicts**: shared memory bank conflicts; high conflicts (>10%) cause serialization; add padding or change access pattern **Compute Profiling:** - **Compute Throughput**: achieved TFLOPS / peak TFLOPS; target 50-80% for compute-bound kernels; A100: 19.5 TFLOPS FP32, 312 TFLOPS FP16 - **SM Efficiency**: percentage of time SMs are active; target 80-100%; low efficiency indicates insufficient work or poor scheduling - **Tensor Core Utilization**: percentage of time Tensor Cores active; target 50-80% for matrix operations; 312 TFLOPS on A100 - **IPC (Instructions Per Cycle)**: instructions executed per cycle; higher is better; target 2-4 for well-optimized kernels **Occupancy Analysis:** - **Achieved Occupancy**: percentage of maximum warps active; target 50-100%; higher occupancy hides latency - **Theoretical Occupancy**: maximum possible based on resource usage; limited by registers, shared memory, block size - **Occupancy Limiter**: identifies limiting factor (registers, shared memory, block size); guides optimization - **Occupancy Calculator**: CUDA Occupancy Calculator spreadsheet; predicts occupancy from resource usage; useful for tuning **Warp Efficiency:** - **Warp Execution Efficiency**: percentage of active threads in executed warps; target >90%; low efficiency indicates divergence - **Branch Efficiency**: percentage of branches without divergence; target >90%; divergent branches cause serialization - **Predication Efficiency**: percentage of instructions not predicated off; target >90%; high predication indicates divergence - **Optimization**: minimize divergence; use ballot/shuffle for divergent code; restructure algorithms **Roofline Model:** - **Concept**: plots achieved performance vs arithmetic intensity; shows whether memory-bound or compute-bound - **Memory Roof**: horizontal line at peak memory bandwidth; memory-bound kernels hit this ceiling - **Compute Roof**: diagonal line at peak compute throughput; compute-bound kernels hit this ceiling - **Optimization**: move toward upper-right (higher intensity, higher performance); tiling increases intensity **Timeline Analysis:** - **Kernel Gaps**: idle time between kernels; indicates launch overhead or synchronization; use streams to overlap - **Memory Transfer Gaps**: idle time during transfers; use async transfers and streams; overlap with compute - **CPU-GPU Sync**: cudaDeviceSynchronize() causes gaps; minimize synchronization; use events for fine-grained control - **Multi-GPU**: visualize cross-GPU communication; identify load imbalance; optimize data distribution **NVTX Markers:** - **Purpose**: annotate code regions; shows in Nsight Systems timeline; helps identify bottlenecks in application logic - **API**: nvtxRangePush("label"), nvtxRangePop(); marks code regions; nvtxMark("event") for single events - **Use Cases**: mark training iterations, data loading, preprocessing, inference; correlate with GPU activity - **Best Practice**: annotate all major code sections; hierarchical markers; color-code by category **Debugging Tools:** - **cuda-memcheck**: detects memory errors; out-of-bounds access, race conditions, uninitialized memory; run with cuda-memcheck ./app - **Compute Sanitizer**: newer tool replacing cuda-memcheck; more features; memcheck, racecheck, initcheck, synccheck modes - **CUDA_LAUNCH_BLOCKING=1**: serializes all operations; easier debugging; disables async; use only for debugging - **cuda-gdb**: command-line debugger; breakpoints, watchpoints, inspect variables; cuda-gdb ./app **Performance Metrics:** - **Achieved Bandwidth**: GB/s of memory traffic; compare to peak (1.5-3 TB/s); target 80-100% for memory-bound - **Achieved TFLOPS**: floating-point operations per second; compare to peak (20-80 TFLOPS); target 50-80% for compute-bound - **Kernel Time**: total kernel execution time; identify slow kernels; focus optimization efforts - **Launch Overhead**: time between kernel launches; target <1% of total time; use CUDA Graphs to reduce **Bottleneck Identification:** - **Memory-Bound**: <50% compute throughput, high memory throughput; optimize memory access patterns, use shared memory, reduce accesses - **Compute-Bound**: <50% memory throughput, high compute throughput; use Tensor Cores, increase ILP, reduce divergence - **Latency-Bound**: low occupancy, low throughput; increase occupancy, reduce register usage, increase block size - **Instruction-Bound**: high instruction overhead; reduce branches, use warp primitives, optimize control flow **Optimization Workflow:** - **Profile**: run Nsight Compute and Nsight Systems; identify bottleneck (memory, compute, latency) - **Analyze**: examine relevant metrics; memory throughput, compute throughput, occupancy, warp efficiency - **Optimize**: apply targeted optimizations; memory coalescing, shared memory, occupancy tuning, divergence reduction - **Measure**: re-profile; verify improvement; compare metrics before and after - **Iterate**: repeat for next bottleneck; diminishing returns after 3-5 iterations; 2-10× total speedup typical **Common Profiling Patterns:** - **Baseline**: profile unoptimized code; establish baseline metrics; identify major bottlenecks - **Incremental**: optimize one aspect at a time; measure impact; easier to attribute improvements - **Comparison**: compare against reference implementation (cuBLAS, cuDNN); identify gaps; target 80-95% of library performance - **Regression**: profile after code changes; detect performance regressions; maintain performance over time **Multi-GPU Profiling:** - **Nsight Systems**: visualizes all GPUs simultaneously; shows cross-GPU communication; identifies load imbalance - **NCCL Profiling**: NCCL_DEBUG=INFO shows communication details; bandwidth, latency, algorithm selection - **Per-GPU Metrics**: profile each GPU separately; identify stragglers; optimize slowest GPU first - **Scaling Analysis**: measure scaling efficiency; compare 1 GPU vs N GPUs; target 80-95% efficiency **Advanced Profiling:** - **Sampling**: sample-based profiling for long-running applications; lower overhead; nsys profile --sample=cpu,cuda - **Metrics Collection**: collect specific metrics; ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed; reduces overhead - **Kernel Replay**: replay kernel with different configurations; find optimal launch parameters; ncu --launch-count 1 --replay-mode kernel - **Source Correlation**: correlate metrics with source code; identify hot spots; ncu --source-level-analysis **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-80 TFLOPS); use Tensor Cores, high ILP, minimal divergence - **Occupancy**: 50-100%; balance register and shared memory usage; 256 threads per block typical - **Warp Efficiency**: >90%; minimize divergence; uniform control flow - **Kernel Time**: <1ms for small kernels, <100ms for large; longer kernels risk timeout; split if necessary **Best Practices:** - **Profile Early**: profile from the start; avoid premature optimization but measure early; establish baseline - **Profile Often**: profile after each optimization; verify improvement; catch regressions - **Use Both Tools**: Nsight Compute for kernel details, Nsight Systems for system view; complementary insights - **Focus on Bottlenecks**: optimize slowest kernels first; 80/20 rule applies; 20% of kernels often account for 80% of time - **Measure, Don't Guess**: intuition often wrong; always measure; data-driven optimization **Common Mistakes:** - **Optimizing Wrong Thing**: optimizing fast kernels instead of slow ones; profile to identify bottlenecks - **Ignoring Occupancy**: assuming higher occupancy always better; balance with resource usage; profile to find optimal - **Over-Optimizing**: diminishing returns after 3-5 iterations; 2-10× total speedup typical; know when to stop - **Not Profiling**: relying on intuition; guessing bottlenecks; always measure actual performance **Real-World Impact:** - **Matrix Multiplication**: profiling reveals 20% of peak; optimization achieves 80-95% of peak; 4-5× speedup - **Reduction**: profiling shows bank conflicts; optimization eliminates conflicts; 2-3× speedup; 60-80% of peak - **Convolution**: profiling reveals memory-bound; shared memory tiling achieves 70-90% of peak; 5-10× speedup - **Custom Kernels**: profiling guides optimization; 2-10× improvement typical; achieves 50-80% of peak GPU Profiling and Debugging represent **the essential tools for GPU performance optimization** — by providing detailed metrics, timeline visualization, and error detection through Nsight Compute and Nsight Systems, developers identify bottlenecks, optimize resource utilization, and achieve 2-10× performance improvement through data-driven optimization, making profiling the difference between GPU code that achieves 10% or 80% of theoretical peak performance where measurement is the only path to understanding actual performance characteristics and intuition often misleads.