Kernel Autotuning

Definition

Kernel autotuning empirically selects a kernel’s launch/compilation parameters — tile/block sizes, num_warps, num_stages, and similar — by benchmarking candidate configurations by wall-clock and caching the fastest, rather than modeling the hardware. The optimal config depends jointly on the problem shape and the GPU, and the search space is discrete and non-obvious, so measurement beats analysis. In Triton this is @triton.autotune; the same idea powers CUTLASS heuristics, cuBLAS kernel dispatch, and torch.compile’s max-autotune.

How it works

Wrap a kernel with a list of triton.Configs and a key; on the first call for each distinct key value, Triton times every config and caches the winner:

@triton.autotune(
    configs=[
        triton.Config({"BLOCK_M":128,"BLOCK_N":256,"BLOCK_K":64}, num_stages=3, num_warps=8),
        triton.Config({"BLOCK_M":64, "BLOCK_N":64, "BLOCK_K":32}, num_stages=5, num_warps=2),
    ],
    key=["M","N","K"],   # re-tune when any of these change
)

The parameters that matter most:

• Block/tile sizes (BLOCK_M/N/K) — set SMEM-tiling and register-blocking extent: bigger tiles → more data reuse and higher arithmetic intensity, but more SMEM/registers and thus lower occupancy (and risk of register spills to local memory). Must respect hardware limits (threads ≤ 1024, SMEM ≤ ~228 KB, dims multiples of 4/16 for vectorized loads / Tensor Cores).
• num_warps — thread-level parallelism per block; num_warps=8 → the tile spreads across $8\times 32=256$ threads.
• num_stages — software-pipelining depth of the K-loop (double/triple buffering of loads via cp.async); most useful for matmul on SM80+; set to 1 to disable when debugging.

Empirical, not modeled. The autotuner just times each config and keeps the fastest — so (1) the first call per new key is slow (keep the config list focused), and (2) every shape parameter that changes the optimal config must be in key, or you silently reuse a stale winner. Set TRITON_PRINT_AUTOTUNING=1 to see the choice. This is the automation of the hand-rolled tile-size sweep from GEMM tuning: the config search grid is discrete precisely because vectorization/MMA alignment and SMEM/register budgets constrain valid shapes.

Why it matters

The sweet spot in the SMEM/register/occupancy trade-off moves with every shape and GPU generation, and no closed-form model reliably finds it — which is why cuBLAS ships hundreds of pre-tuned kernels and CUTLASS/Triton lean on autotuning. For a frontier lab running one model shape across a fleet, a good autotuned config is a direct throughput (and cost) multiplier, and re-tuning is mandatory when moving between A100/H100/B200 or changing sequence length, batch, or precision.

Taught in

• 07-triton-and-modern-kernels — §7 @triton.autotune, the key knobs, and its empirical cost.
• 05-optimizing-gemm — the hand-rolled tile-size sweep autotuning automates.

See also

• triton — where @triton.autotune lives
• torch-compile — autotunes generated kernels (max-autotune)
• shared-memory-tiling · register-blocking — what block sizes control
• occupancy — the trade-off being tuned
• arithmetic-intensity · gemm · tensor-core
• gpu-systems-for-llms