analyze-kernel-bottleneck

Files (1)

• SKILL.md

SKILL.md (first lines)

# 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:
   ```bash
   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:
   ```bash
   ./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`

**Got:** Baseline numbers: kernel time in ms, effective GFLOPS, and effective bandwidth.

**If fail:** Check that the kernel launches without error (`CHECK_CU` macro). Verify warmup runs precede measurement. Ensure problem dimensions are large enough to saturate the GPU (small problems may bottleneck on launch overhead).

### Step 2: Classify on the Roofline

Compute arithmetic intensity and compare against the machine balance point to classify the kernel:

1. **Calculate arithmetic intensity**: `AI = FLOPs / bytes_loaded_from_global_memory`. Count only unique bytes loaded from DRAM (not shared memory or register reuse).
2. **Look up machine balance point**: `balance = peak_compute / peak_bandwidth`.
3. **Classify**: If `AI < balance`, the kernel is memory-bound. If `AI > balance`, the kernel is compute-bound.

**GA104 (RTX 3070 Ti) Reference Values:**

| Resource | Peak | Unit |
|---|---|---|
| FP32 FFMA | 21.7 | TFLOPS |
| FP16 Tensor Core (HMMA) | 174 | TFLOPS |
| INT8 Tensor Core (IMMA) | 696 | TOPS |
| DRAM Bandwidth | 608 | GB/s |
| L2 Cache | 4 | MB |
| SMs | 48 | |

**Derived Balance Points:**

| Precision | Balance Point (FLOP/byte) |
|---|---|
| FP32 FFMA | 21700 / 608 = 35.7 |
| FP16 TC | 174000 / 608 = 286.2 |
| INT8 TC | 696000 / 608 = 1144.7 |

4. **Compute attained fraction**: `attained = effective_throughput / peak_throughput`. If memory-bound: compare effective bandwidth to 608 GB/s. If compute-bound: compare effective GFLOPS to the relevant peak.

**Got:** Classification as compute-bound, memory-bound, or latency-bound (low occupancy causing neither compute nor memory saturation) with numerical justification.

**If fail:** Recheck byte counting. Watch for redundant re-reads (e.g., 9x in direct conv2d without im2col). If neither compute nor memory is saturated, the kernel is likely latency-bound (see Step 3).

### Step 3: Calculate Occupancy

Determine active warps per SM from the launch configuration and resource usage:

1. **Extract resource usage**:
   ```bash
   nvcc --cubin -arch=sm_86 -O2 --resource-usage -o kernel.sm_86.cubin kernel.cu 2>&1 | grep -E 'registers|smem'
   ```
2. **From launch config**: `warps_per_block = threads_per_block / 32`.
3. **Compute blocks/SM** from each limiting factor:
   - Register limit: `floor(65536 / (registers_per_thread * threads_per_block))`
   - Smem limit: `floor(available_smem_per_SM / smem_per_block)` -- see Step 6 for cliff
   - Warp limit: `floor(48 / warps_per_block)` (GA104 max: 48 warps/SM)
   - Block limit: 16 blocks/SM max on GA104
4. **Actual blocks/SM** = `min(register_limit, smem_limit, warp_limit, block_limit)`.
5. **Active warps/SM** = `blocks_per_SM * warps_per_block`.
6. **Key threshold**: 8 warps/SM is sufficient for latency hiding on GA104. Below 8 = structural problem causing latency-bound behavior.

**Got:** Occupancy table showing blocks/SM, active warps/SM, and the limiting factor (registers, smem, or warps).

**If fail:** Check `cuFuncSetAttribute` for dynamic shared memory. Verify `--resource-usage` reports match the actual launch configuration. If register count is unexpectedly high, try `--maxrregcount=N` to cap registers (trading register spills for occupancy).

### Step 4: Compute the Compute/Load Ratio Per Tile

Count compute instructions and load bytes per K-tile from SASS (not source code):

1. **Disassemble**:
   ```bash
   cuobjdump -sass kernel.sm_86.cubin > kernel.sass
   ```
2. **Count compute instructions per tile** (the inner loop over one K-tile):
   - `grep -c 'HMMA' kernel.sass` -- FP16 Tensor Core ops
   - `grep -c 'IMMA' kernel.sass` -- INT8 Tensor Core ops
   - `grep -c 'FFMA' kernel.sass` -- FP32 fused multiply-add
3. **Count global loads per tile**:
   - `grep -c 'LDG' kernel.sass` -- global memory loads
   - Multiply by bytes per load (typically 16 bytes for LDG.128)
4. **Calculate ratio**: `compute_ops / load_ops` per tile.
5. **Classify** using the cp.async decision threshold (from gpu_reflections.md Insight 2):
   - **High** (>20:1): cp.async is net-negative; warp interleaving already hides DRAM latency. Focus on algorithmic changes. Reference: Flash Attention has 64 HMMA per tile = high ratio, cp.async measured -5%.
   - **Medium** (5-20:1): cp.async may help, benchmark both paths.

From pjt222/agent-almanac, MIT.

More from pjt222/agent-almanac