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.**