Autotuning

01 · Why guessing loses

The best block shape isn't a constant — it's a function of your GPU

Day 14 recommended 256 threads per block and day 21 picked BLOCK=4096 on faith. Both numbers are actually functions — of shared memory per SM, register file size (day 09), warp count (day 06), L1 geometry (day 10). Change the GPU and the optimum moves: the block shape that hits 90% of bandwidth on a 4070 can leave 30% on the table on an H100, and vice versa. Hand-tuning per GPU doesn't scale past your own desk — and open-source kernels can't hand-tune for hardware they've never seen.

Triton's answer is refreshingly blunt: if the parameter matters, measure it. Declare a small set of candidate configurations; on first launch, run each and keep the fastest; cache the winner keyed on the arguments you nominate.

02 · The decorator

@triton.autotune, end to end

@triton.autotune(
    configs=[
        triton.Config({'BLOCK': 256}, num_warps=4),
        triton.Config({'BLOCK': 1024}, num_warps=4),
        triton.Config({'BLOCK': 4096}, num_warps=8),
    ],
    key=['n'],   # re-tune when the problem size changes
    reset_to_zero=['out_ptr'],  # for kernels that accumulate
)
@triton.jit
def add(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < n
    tl.store(out_ptr + offs,
              tl.load(x_ptr + offs, mask=mask)
               + tl.load(y_ptr + offs, mask=mask),
              mask=mask)

The key list is the part worth staring at. Tuning results are cached per distinct value of the keyed arguments, because the best BLOCK for n = 1,000 is not the best for n = 100,000,000 — different shapes mean different grids, different tail-masking costs, different occupancy. Key on whatever changes the optimal shape; production kernels typically key on tensor shapes, never on raw pointers or batch contents. The first call for each new shape pays the search; every call after is a dictionary lookup.

03 · What's actually being searched

Three knobs, all familiar under other names

Notice the interactions you already understand: a bigger block wants more shared memory (day 08) and more registers (day 09), which cuts programs-per-SM (day 06), which changes latency hiding (day 01). Deeper pipelining hides memory latency but holds more tiles live. No closed-form model survives all three interacting — which is precisely why search beats theory here, and why day 14 said "default to 256, then measure": autotune is the measuring, industrialized.

04 · Costs, caches, and production practice

Tune once at startup, not inside the hot loop

The search runs real kernel launches with timing around them — for a big GEMM with a wide config list, that's seconds of one-time cost. Three practices keep it sane. Budget the warmup: tune during model loading, before latency matters. Prune the list: eight sensible configs beat fifty random ones. Pre-bake when it matters: inference engines like vLLM snapshot tuned configs per GPU and skip the search entirely — deterministic startup, no surprise timings on the first user request. And never put an autotuned kernel's first call inside a benchmark: you'll measure the search, not the kernel.

Next: day 23 — kernel fusion: why frameworks fuse elementwise ops, what it saves in memory traffic, and a calculator to watch the savings stack up.