GPU Profiling Nsight Compute

GPU Profiling with Nsight Compute is the systematic analysis of GPU kernel performance characteristics — including compute throughput, memory throughput, occupancy, stall reasons, and instruction mix — to identify bottlenecks and guide optimization decisions using the detailed hardware performance counters available on NVIDIA GPUs.

Key Profiling Metrics:
- SM Throughput: percentage of peak compute throughput achieved — low values indicate instruction-level inefficiency (poor ILP, warp divergence, or stalls)
- Memory Throughput: percentage of peak memory bandwidth utilized — high values indicate a memory-bound kernel; optimization should focus on reducing memory traffic or improving access patterns
- Occupancy: ratio of active warps to maximum warps per SM — higher occupancy helps hide latency but isn't always necessary; some kernels achieve peak performance at 50% occupancy with good data reuse
- Warp Execution Efficiency: average number of active threads per warp instruction — values below 32 indicate divergence; target >28 for well-optimized kernels

Stall Analysis:
- Memory Dependency Stalls: warps waiting for memory load/store completion — indicates insufficient occupancy or poor memory access patterns (uncoalesced, cache misses)
- Execution Dependency Stalls: warps waiting for previous instruction result — indicates long instruction latency chains (transcendental functions, integer division) without sufficient parallelism to hide latency
- Synchronization Stalls: warps waiting at __syncthreads() or atomics — indicates load imbalance within a block or excessive atomic contention
- Instruction Fetch Stalls: instruction cache misses, typically from very large kernels or low I-cache locality — rare but occurs with complex control flow and large instruction footprints

Memory Analysis:
- L1/L2 Cache Hit Rate: percentage of loads served from cache vs. DRAM — low hit rates suggest poor data locality or working set larger than cache capacity
- Sector Utilization: percentage of bytes in each cache sector (32 bytes) actually used by the requesting warp — low utilization indicates poor coalescing or wasted bandwidth from partial cache line usage
- Shared Memory Efficiency: transactions per request — 1.0 means no bank conflicts; higher values indicate N-way conflicts reducing shared memory bandwidth
- DRAM Read/Write Ratio: excessive writes relative to reads may indicate unnecessary store operations or write-back traffic — read-heavy workloads are more common for inference-style kernels

GPU profiling with Nsight Compute is the indispensable diagnostic tool for GPU performance engineering — without quantitative profiling data, kernel optimization is guesswork; with it, engineers can systematically identify and eliminate bottlenecks to approach the theoretical performance ceiling defined by the roofline model.