| name | tuning-triton |
| description | Tune Triton and Gluon GPU kernels on AMD Instinct — build a config space including the AMD-only knobs (waves_per_eu, matrix_instr_nonkdim, kpack), prune it before racing, set the autotune key correctly, and confirm the win exceeds the noise floor. Use when the kernel you are tuning is written in Triton or Gluon. |
Tuning Triton kernels
Read ../tuning-core/SKILL.md first — this skill is the Triton specialization of that loop.
Triton is the one backend where you tune by authoring the search space, not by invoking a
tuner. @triton.autotune will faithfully race whatever list you hand it, so the quality of
your result is capped by the quality of that list. Three things decide the outcome:
- whether the AMD-specific knobs are in the space at all (they are easy to omit silently);
- whether the
key= is right (getting it wrong costs multiples, silently);
- whether you pruned before racing (an unpruned space wastes most of the compile budget).
Measurements below are gfx942 / MI300X, Triton 3.6.0, bf16.
1. Put the AMD knobs in the space
This is the trap that costs the most performance, because it fails as a TypeError you can
"fix" by deleting the knob.
triton.Config.__init__ accepts only num_warps, num_stages, num_ctas, maxnreg,
pre_hook, ir_override. Passing an AMD knob as a keyword raises:
triton.Config({'BLOCK_M': 64}, waves_per_eu=2)
The knobs are real, but they travel in the first positional dict, alongside your
tl.constexpr block sizes:
triton.Config({'BLOCK_M': 64, 'waves_per_eu': 2, 'matrix_instr_nonkdim': 16, 'kpack': 2},
num_warps=8, num_stages=2)
Verify rather than trust: cfg.all_kwargs() should list them. They are consumed by the HIP
backend, whose full option set you can read on any box with
import dataclasses
from triton.backends.amd.compiler import HIPOptions
[f.name for f in dataclasses.fields(HIPOptions)]
The three worth searching, and what they trade:
| knob | default | what it controls | when it matters |
|---|
matrix_instr_nonkdim | 0 (auto) | MFMA instruction shape (e.g. 16 → 16×16, 32 → 32×32) | most impactful; small tiles and skinny shapes often want 16 |
kpack | 1 | K-elements packed per load into the MFMA operand | second-order; interacts with BLOCK_K |
waves_per_eu | 0 (auto) | occupancy hint — caps registers to fit more waves per EU | latency-bound kernels; can hurt register-heavy ones |
Measured on 4096³ bf16, tile 128×128×64, num_warps=8, num_stages=2, on both parts
(tuning_benchmark/tools/triton_tables.py reproduces this):
| variant | gfx942 TFLOPS | spread | vs base | gfx950 TFLOPS | spread | vs base |
|---|
| baseline (all auto) | 356.8 | 5.4% | — | 760.5 | 5.1% | — |
matrix_instr_nonkdim=16 | 384.3 | 0.8% | +7.7% | 791.1 | 1.6% | +4.0% |
matrix_instr_nonkdim=32 | 372.5 | 7.7% | +4.4% | 757.9 | 2.8% | −0.3% |
Read that table the way the next section says to, not the way it first looks. The conclusion
survives the move — nonkdim=16 is the real win on both parts, nonkdim=32 is not
distinguishable from nothing on either — but the magnitude halves and nonkdim=32 crosses
from marginally positive to marginally negative. That is the general pattern for this whole
skill on gfx950: the knob that matters still matters, and by less.
The gfx950 column was itself only measurable after interleaving the three variants. Timed
back-to-back on that part they report 13.8–20.7% spreads, which swallows the entire result and
would have made this table read as three indistinguishable numbers. See
../tuning-core/measurement.md Rule 6b before reproducing any of this on MI355.
2. Compare the win to the noise floor, not to zero
Both variants above are "faster than baseline". Only one of them is a result.
nonkdim=16: +7.7%, and its own spread is 0.8%. The effect is an order of magnitude
larger than the measurement's own variability. Real.nonkdim=32: +4.4%, spread 7.7%. The claimed gain is smaller than the noise in the
measurement that produced it. Not distinguishable from nothing.
A single do_bench call reports neither of these facts. It returns one number for each and
they both look like wins. The discipline — take independent repeats, report the median and
the spread, and refuse to believe a delta smaller than the spread — is what separates them.
Full rules in ../tuning-core/measurement.md.
Practical consequence for autotuning: @triton.autotune picks the single fastest config from
one timing pass. When several configs sit inside each other's noise band, which one "wins" is
partly chance. That is acceptable — you get one of the good configs — but it means you must
not report the autotuner's margin as a speedup without re-measuring the winner properly.
3. Set key= to every variable that changes the answer
key is the list of kernel arguments whose values invalidate the cached choice. Autotuning
runs once per distinct key tuple; every later call with a matching tuple reuses that config.
Omit a dimension and you silently serve a config tuned for a different shape. Same kernel,
same config list, only key differs:
| shape | key=['M','N','K'] | key=['N','K'] |
|---|
| M=4096 N=4096 K=4096 | 348.9 µs — picks BM=256,BN=128 | 350.2 µs — picks BM=256,BN=128 |
| M=1 N=4096 K=4096 | 43.4 µs — picks BM=16,BN=64 | 162.9 µs — reuses BM=256,BN=128 |
3.75× slower on the decode shape, no error, no warning. The big-M config was cached
first; M is not in the key, so the M=1 call is treated as the same problem. A BM=256 tile
on a 1-row GEMM computes 255 rows of padding.
On gfx950 the same experiment gives 2.45× (M=1: 54.5 µs with the decode tile, 22.2 µs with
the right one — and 180.6 µs vs 479.7 µs at M=4096 in the other direction). Smaller, still
multiples, still silent. The penalty shrinks because a 256-row tile wastes proportionally less
on a part with more bandwidth per CU, not because the mistake became safe.
Rules that follow:
- include every dimension the kernel branches or tiles on — for a GEMM that is
M, N, K;
- include dtype-selecting or feature-flag arguments if they change which code path runs;
- if a value is genuinely constant, make it a
tl.constexpr rather than trusting yourself to
remember it is constant.
Cost of being right: one autotune pass per distinct shape. In serving, bound that by
bucketing M (see ../tuning-core/search_strategy.md) rather than by shortening the key.
4. Prune before you race
Every config in the list costs a compile plus a timing pass, and the space is multiplicative:
block sizes × warps × stages × three AMD knobs reaches thousands of entries. Most are
invalid — they exceed LDS, or give threads no work — and Triton finds that out by compiling
them.
Filter analytically first. The pruning predicate and its measured reduction (1600 → 912
configs, 57%) are in ../tuning-core/search_strategy.md; example_autotune.py here shows it
wired into a real kernel.
Then narrow by regime — this is worth more than it sounds. Same kernel, same autotuner, same
M=1 N=4096 K=4096 shape, only the space differs:
| space raced | winner | gfx942 | gfx950 |
|---|
general (BM ∈ 64,128,256) | BM=64, BN=64, BK=64, nonkdim=16 | 0.041 ms | 0.041 ms |
decode (BM ∈ 16,32, deep BK) | BM=16, BN=32, BK=256, nonkdim=0, warps=2 | 0.027 ms | 0.023 ms |
34% faster from changing the candidate list alone on gfx942, and 44% on gfx950. The general
space simply did not contain the right answer, and an autotuner cannot select what you did not
offer it. Note the winner also chose nonkdim=0 — the auto MFMA shape — where the square case
wanted 16; another reason a single fused list serves neither regime well.
This is the one measurement in this skill that got larger on the newer part. Splitting the
space by regime is the highest-value item here on gfx950, above any individual knob.
One caveat the example makes explicit: @triton.autotune is evaluated at import time, so the
config list is fixed before any argument is seen. You cannot choose the space from the shape
inside the decorator. Either accept one larger union list, or define separate kernels per
regime and dispatch between them yourself — a real cost of regime-splitting worth pricing in.
You can also use Triton's own hook, prune_configs_by=, which additionally lets you drop
configs using the runtime argument values (early_config_prune, perf_model, top_k).
5. Control the cache while you tune
Autotune results and compiled binaries are cached under /root/.triton/cache
(knobs.cache.dir). During a tuning session this is a hazard: you change the config list,
rerun, and get the previous answer.
export TRITON_PRINT_AUTOTUNING=1
rm -rf /root/.triton/cache
TRITON_PRINT_AUTOTUNING=1 is the cheapest engagement signal Triton offers: it tells you
autotuning actually ran and which config won. If it prints nothing, you are on a cached
result and your edits did not take effect. Also readable in-process as
kernel.best_config after the first call — that is how the key= table above was produced.
6. Gluon
Gluon (triton.experimental.gluon) is the lower-level dialect: you place data layouts by
hand instead of letting the compiler infer them. Everything above still applies — same
Config, same key, same noise discipline — but the search space gains layout choices
(BlockedLayout, SwizzledSharedLayout, PaddedSharedLayout, DotOperandLayout,
AMDMFMALayout).
The arch split is explicit in the API, which makes it a good place to reason about gfx942 vs
gfx950:
from triton.experimental.gluon.language.amd import cdna3, cdna4
cdna3 (gfx942): mfma, buffer_load/buffer_store, buffer_atomic_*cdna4 (gfx950): all of the above plus mfma_scaled, get_mfma_scale_layout,
async_copy
So a Gluon kernel written against cdna3 runs on both; one written against cdna4 does not
run on gfx942. Tune on the arch you will deploy on — and note that mfma_scaled is the entry
point for the microscaled dtypes that only gfx950 has, which is why those shapes have no
gfx942 equivalent to carry a tuned config over from.
Reach for Gluon only after the Triton-level space is exhausted: the extra layout freedom is
also extra ways to be slow, and the compiler's inferred layouts are usually good.
Tuning someone else's Gluon kernel: the space is smaller than the signature
A hand-placed layout can encode the tile shape, and then the tile shape stops being tunable.
aiter's gluon.gemm_afp4wfp4 accepts the usual ten-key GEMM config and writes its layouts as
literal basis vectors:
linear_mn = gl.DistributedLinearLayout(
reg_bases=[[0, 1], [0, 2], [0, 4], [0, 16], [0, 128], [64, 0], [128, 0]],
warp_bases=[[0, 32], [0, 64], [32, 0]],
shape=[BLOCK_SIZE_M, BLOCK_SIZE_N],
)
A basis of 128 needs a dimension above 128, and three warp bases mean eight warps. The tile
is pinned at 256×256×256 on num_warps=8 and nothing in the signature says so. Measured on
gfx950, one knob moved at a time:
| knob | result |
|---|
BLOCK_SIZE_M/N/K, num_warps | LLVM ERROR ... abort() — the process dies |
num_stages, matrix_instr_nonkdim | accepted, declared, never read; median moves <2% |
GROUP_SIZE_M, waves_per_eu, NUM_KSPLIT, cache_modifier | real |
Three consequences worth carrying to any hand-written-layout kernel:
Probe the space in subprocesses. The tile-shape failure is abort(), not an exception —
no try will catch it, and an in-process sweep loses every result it had gathered. One
subprocess per candidate, judged by exit code, is the only safe way to map the boundary. Do
it once and turn the answer into a constraint.
A thin config file is evidence about degrees of freedom, not effort. The shipped Gluon
config is a single "any" entry, against M-binned entries across forty-odd N/K-specialised
files for the Triton path to the same math. That looks like an untuned kernel and is not: one
entry covers a four-knob surface. Sweeping it found exactly one win, NUM_KSPLIT=2 at
M=16 (11.2%), which is the axis you would predict matters at decode shapes and the one thing
a single shape-independent entry cannot get right.
Declared is not read. num_stages and matrix_instr_nonkdim are in the kernel signature
and absent from the body — Gluon pipelines by hand and the MFMA shape is hardcoded. Leaving
them in the tunable list buys recompiles of identical machine code, and their apparent 1-2%
"wins" are noise that a budget-limited sweep will happily report.
7. Then verify engagement
A tuned Triton config is only useful if the framework calls your kernel. In vLLM and SGLang
many Triton kernels are selected by env var or by a config JSON keyed on shape and device
name — and the device-name lookup has its own failure mode (see ../tuning-in-vllm/). Finish
with ../tuning-core/engagement_verification.md.
Files
example_autotune.py — a GEMM with pruning, AMD knobs, correct key, and a noise-aware
comparison. Runnable; written to be read and adapted, not imported.
Checklist
