CUDA Optimization Patterns: Kernel Fusion, Shared Memory Tiling, and Occupancy Tuning for Transformer Inference and Training

The FlashAttention series exemplifies this hierarchy: FlashAttention 2 achieves 2-3x speedup primarily by reducing HBM reads/writes through tiling to shared memory, not by reducing FLOPs. FlashAttention 3 adds asynchronous processing using CUDA warp group matrix multiply-and-accumulate (WGMMA) instructions on Hopper.

Understanding where to invest engineering time matters: optimizing memory-bound kernels can yield more benefit than optimizing compute-bound ones.

Kernel fusion reduces launch overhead and avoids intermediate HBM writes. For transformer inference, fusing layer normalization, residual add, attention, and feed-forward eliminates 3-4 intermediate HBM round-trips per layer. NVIDIA TensorRT-LLM and ThunderKittens provide fused kernels for standard patterns.

CUDA Graphs reduce launch overhead for a 96-layer Llama inference from 850 µs to 45 µs per decode step, a 19x reduction on H100.

Shared memory tiling splits Q, K, V matrices into tiles that fit in shared memory (228 KB per SM on H100, 256 KB on B200). The online softmax algorithm recomputes normalization over the tiled KV without writing intermediate attention scores to HBM.

Optimal tile sizes for H100: Br=64, Bc=128 for FP16 at 4K sequence length. For B200 with 160 SMs: Br=64, Bc=128 or Br=128, Bc=64 depending on sequence length.

CUDA occupancy should target the point where memory latency is fully hidden without register spilling. On H100 with 64 warps per SM, 228 KB shared memory, and 65,536 registers, pushing register usage beyond limits forces spills that degrade performance by 15-30%.

Dynamic shared memory via cudaFuncSetAttribute enables flexible allocation per SM, critical for inference serving with unpredictable request sizes.