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

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

On failure: 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.

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

On failure: 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:

   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.

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

On failure: 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:

   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.
   • Low (<5:1): cp.async strongly beneficial; loads dominate and async copy hides latency. Reference: IGEMM has 8 IMMA per tile = low ratio, cp.async measured +35%.

Expected: Compute/load ratio with classification (high/medium/low) and cp.async recommendation.

On failure: Count from SASS disassembly, not source code -- the compiler may fuse, eliminate, or reorder instructions. Ensure you are counting instructions within the inner loop only (the K-tile iteration), not the entire kernel.

Step 5: Inspect SASS Instructions

Examine the full SASS instruction mix and stall codes:

1. Disassemble (if not done in Step 4):

   cuobjdump -sass kernel.sm_86.cubin > kernel.sass
   

2. Count key instruction types:

   grep -c 'HMMA.16816' kernel.sass      # FP16 Tensor Core
   grep -c 'IMMA.16816' kernel.sass      # INT8 Tensor Core
   grep -c 'FFMA' kernel.sass            # FP32 fused multiply-add
   grep -c 'LDGSTS' kernel.sass          # cp.async (global->shared)
   grep -c 'LDG' kernel.sass             # Global load
   grep -c 'STS' kernel.sass             # Shared store
   grep -c 'LDS' kernel.sass             # Shared load
   grep -c 'BAR.SYNC' kernel.sass        # Barrier synchronization
   grep -c 'SHFL' kernel.sass            # Warp shuffle (reductions)
   grep -c 'MUFU' kernel.sass            # Special function unit
   

3. Check stall codes on critical instructions:

   grep 'HMMA' kernel.sass | head -5     # Expect S08 minimum (hardware constraint)
   grep 'IMMA' kernel.sass | head -5     # Compiler emits S04, reducible to S02 via CuAssembler
   grep 'FFMA' kernel.sass | head -5     # Check for S04 (reducible to S01 on independent FFMAs)
   

4. Identify optimization targets:
   • HMMA S08 stalls: hardware minimum on Ampere, cannot be reduced. Focus elsewhere.
   • IMMA S04 stalls: compiler is conservative. CuAssembler can tighten to S02 (measured 15-20% gain).
   • FFMA S04 stalls: if independent, reducible to S01 via CuAssembler.
   • Excessive BAR.SYNC: may indicate over-synchronization between pipeline stages.

Expected: Instruction count table and stall code summary with identified optimization targets.

On failure: Ensure cuobjdump architecture matches the kernel compilation target (both must be sm_86). If SASS output is empty, the cubin may be corrupt -- recompile.

Step 6: Check the Smem Cliff

Determine whether shared memory usage crosses the architecture-specific occupancy cliff:

1. Read smem/block from --resource-usage output (Step 3) or cuobjdump --res-usage kernel.sm_86.cubin.
2. Compare against cliff threshold:
   • GA104 (sm_86): 100 KB max smem/SM. Cliff at 50 KB/block.
   • Confirmed empirically: 48 KB/block -> 2 blocks/SM (good), 56 KB/block -> 1 block/SM (2x regression).
3. If above cliff (smem > 50 KB/block):
   • Blocks/SM drops to 1, active warps drop to warps_per_block (typically 4).
   • 2x performance regression expected from exposed DRAM stalls.
4. Check double-buffering impact: Double-buffering doubles smem usage. If current smem is 30 KB, double-buffered = 60 KB, which crosses the cliff. Evaluate whether the async benefit outweighs the occupancy loss.
5. Record smem/block, blocks/SM, and whether the cliff is crossed.

Expected: Smem/block value with blocks/SM count and explicit statement of whether the 50 KB cliff is crossed.

On failure: If above cliff and occupancy is the bottleneck, the optimization strategy must change: reduce tile size to get smem under 50 KB, or accept

---

Content truncated.