gpu kernel profiling
**GPU Kernel Profiling** is the **systematic measurement and analysis of GPU kernel execution characteristics — occupancy, memory throughput, compute utilization, stall reasons, and instruction mix — using profiling tools to identify performance bottlenecks** and guide optimization toward the specific limiter (compute-bound, memory-bound, or latency-bound) that determines kernel performance.
Without profiling, GPU optimization is guesswork. A kernel running at 5% of peak FLOPS might be memory-bound (and unreachable by compute optimization) or might have poor occupancy (fixable by reducing register usage). Profiling reveals which optimization will actually improve performance.
**Profiling Tools**:
| Tool | Vendor | Capabilities |
|------|--------|-------------|
| **Nsight Compute** | NVIDIA | Kernel-level metrics, roofline, source correlation |
| **Nsight Systems** | NVIDIA | Timeline, API trace, CPU-GPU interaction |
| **ROCprofiler** | AMD | Kernel metrics for CDNA/RDNA GPUs |
| **Omniperf** | AMD | High-level performance analysis |
| **Intel VTune** | Intel | GPU profiling for Intel GPUs |
**Key Metrics**:
1. **Occupancy**: Active warps / maximum warps per SM. Low occupancy (<50%) means insufficient parallelism to hide memory latency. Caused by: excessive register usage, excessive shared memory per block, or too-small block sizes. **Achieved occupancy** (runtime average) matters more than theoretical occupancy.
2. **Memory throughput**: Actual bytes/second to/from each memory level vs. peak. Global memory throughput near peak (80%+) with low compute utilization → memory-bound kernel. Shared memory throughput near peak with bank conflict stalls → shared memory optimization needed.
3. **Compute throughput**: Actual FLOP/s vs. peak. Low compute throughput with low memory throughput → latency-bound (insufficient occupancy or instruction-level parallelism).
4. **Warp stall reasons**: Nsight Compute breaks down why warps are stalled: memory dependency (waiting for load), execution dependency (waiting for ALU result), synchronization (`__syncthreads()` barrier), and instruction fetch (instruction cache miss). This directly identifies the bottleneck.
**Roofline Analysis**: The roofline model plots kernel performance (FLOP/s) against arithmetic intensity (FLOP/byte of memory traffic). Kernels below the roofline have optimization opportunity. Memory-bound kernels (left of the ridge point) benefit from reducing memory traffic (tiling, caching, compression). Compute-bound kernels (right of the ridge point) benefit from algorithmic optimization or mixed-precision arithmetic.
**Profiling Methodology**: 1) Profile baseline kernel with Nsight Compute. 2) Identify primary bottleneck (memory, compute, latency). 3) Apply targeted optimization (not random optimization). 4) Re-profile to verify improvement and identify next bottleneck. 5) Iterate until satisfied. Each optimization typically shifts the bottleneck to a different resource — the art is knowing when the kernel is "close enough" to the hardware limit.
**GPU kernel profiling transforms performance optimization from art to science — by quantifying exactly where execution time is spent and why, profiling enables targeted optimizations that deliver measurable improvement rather than hopeful speculation, making it the indispensable first step in any GPU optimization effort.**