analyze-kernel-bottleneck

Analyze Kernel Bottleneck

Systematically identify whether a GPU kernel is compute-bound, memory-bound, or latency-bound by measuring baseline performance, classifying on the roofline, computing occupancy and compute/load ratio per tile, inspecting SASS instruction mix and stall codes, checking the shared memory cliff, and applying a decision matrix to select the right optimization strategy.

When to Use

• Before optimizing any CUDA kernel -- establish baseline and classify bottleneck type
• After writing a first working version of a kernel to identify the optimization path
• When a kernel underperforms expectations relative to theoretical peak
• When deciding between cp.async, larger tiles, or algorithmic restructuring

Inputs

• Required: Compiled kernel (.cubin or .cu source with build command)
• Required: Benchmark harness that launches the kernel with CUDA event timing
• Required: Problem dimensions (e.g., M, N, K for GEMM; seq_len, heads, head_dim for attention)
• Optional: Target GPU architecture (default: GA104 / sm_86 / RTX 3070 Ti)
• Optional: Expected peak utilization percentage for comparison
• Optional: Prior profiling data (Nsight Compute reports)

Procedure

Step 1: Measure Baseline Performance

Run the kernel with CUDA events (BenchTimer), record time in milliseconds. Calculate effective throughput metrics:

1. Compile the kernel if not already built:

   nvcc --cubin -arch=sm_86 -O2 -o kernel.sm_86.cubin kernel.cu
   nvcc -arch=sm_86 -O2 -o bench bench.cu -lcuda -I../../phase2/common
   

2. Run with representative problem sizes, ensuring warmup runs precede measurement:

   ./bench 4096 4096 4096
   

3. Record kernel time in ms from CUDA events (not wall-clock).
4. Calculate effective GFLOPS and effective bandwidth:
   • GEMM: effective_gflops = (2 * M * N * K) / (time_ms / 1000) / 1e9
   • Bandwidth-limited kernels: effective_bw = total_bytes / (time_ms / 1000) / 1e9
   • Flash Attention: effective_gflops = (4 * batch * heads * seq_len^2 * head_dim) / (time_ms / 1000) / 1e9