Tune GEMM and fused ops through aiter on AMD Instinct — run the gradlib tuner across hipblaslt/triton/flydsl/asm/torch backends, write tuned config CSVs to the directory the runtime actually reads, and prove engagement with AITER_LOG_TUNED_CONFIG. Use when tuning ops dispatched through aiter in vLLM or SGLang.
Tune GEMM and fused ops through aiter on AMD Instinct — run the gradlib tuner across hipblaslt/triton/flydsl/asm/torch backends, write tuned config CSVs to the directory the runtime actually reads, and prove engagement with AITER_LOG_TUNED_CONFIG. Use when tuning ops dispatched through aiter in vLLM or SGLang.
Tuning through aiter
Read ../tuning-core/SKILL.md first.
aiter is where most serving-path GEMMs actually land, and it is the one backend that tunes
across languages: a single gradlib run races hipBLASLt solutions, Triton kernels, FlyDSL
kernels, hand-written assembly and torch against each other, then records the winner per
shape. That makes it the highest-leverage tuning target and the one with the most ways to
produce an artifact that never gets read.
Everything below was measured on gfx942 / MI300X, aiter 0.1.13, in the vllm image.
1. The artifact is a CSV row, keyed on your hardware
Tuned configs are CSV files under aiter/configs/. The bf16 GEMM schema:
The first two columns are the whole story of why tuned artifacts do not travel. A lookup
matches on gfxandcu_num. Check what you are:
from aiter.jit.utils.chip_info import get_cu_num, get_gfx
get_gfx(), get_cu_num() # -> ('gfx942', 304) on MI300X
Then check what shipped. In the vllm image, measured:
config file
rows
gfx/cu_num present
bf16_tuned_gemm.csv
780
gfx950/256 only
a4w4_blockscale_tuned_gemm.csv
1500
gfx950/256 only
a8w8_blockscale_bpreshuffle_tuned_gemm.csv
1498
gfx950/256 only
a8w8_blockscale_tuned_gemm.csv
15611
gfx942/304, gfx950/256
a8w8_bpreshuffle_tuned_gemm.csv
2084
gfx942/304, gfx942/80, gfx950/256
On an MI300X box: 780 tuned bf16 rows, none of them usable. Not a bug — the image ships
for several targets and the bf16 table happens to be populated for gfx950 only. But anyone
assuming "aiter ships tuned configs, so bf16 GEMM is tuned" is wrong on that box, and nothing
says so at runtime unless you ask. cu_num=80 rows in the a8w8 file are from a partitioned
MI300 mode — same gfx, different CU count, still no match at 304.
Run the same image on an MI355X and every one of those rows hits. That is the useful thing
about this table: it is the same artifact, and which half of it is dead depends entirely on
what you plugged in. Two consequences that showed up when the corpus was swept on both parts:
Tuning headroom is a property of the box, not of the library. The same sweep over the
same cases returned a dispatch-weighted 1.96x on gfx942 and 1.40x on gfx950, and the gfx950
baselines were already 2-3.6x faster in absolute terms. Less headroom because more of the
shipped tables applied, not because the tuner got worse. Reporting an uplift ratio without
saying which tables were live is reporting an artifact of the image.
A well-populated table can still ship an unsound row. On gfx950, the shipped
gemm_a16w8_blockscale config selects NUM_KSPLIT=8 for M≤16 at K=2880 — but K/group_k is
22.5 groups, so the split boundaries cut quantization groups and the baseline is simply
wrong, by ~40%, before any tuning starts. A harness that trusts the shipped config as its
reference measures every candidate against a wrong answer. Check the baseline's correctness
first; see ../tuning-core/correctness_gates.md.
This is the concrete form of the general rule: a tuned artifact is never portable across
architectures. Re-tune per target — and verify the target's own rows before trusting them.
2. Turn on the engagement signal before you change anything
export AITER_LOG_TUNED_CONFIG=1
aiter then prints, per call, whether the lookup hit. Untuned:
[aiter] shape is M:4096, N:4096, K:4096 dtype='torch.bfloat16' ... not found tuned config
in /tmp/aiter_configs/bf16_tuned_gemm.csv, will use default config! using torch solution:0
Tuned:
[aiter] shape is M:4096, N:4096, K:4096 ... found padded_M: 4096, N:4096, K:4096
is tuned on cu_num = 304 in /tmp/aiter_configs/bf16_tuned_gemm.csv, libtype is hipblaslt
Take the baseline reading first. If the shape already says is tuned on cu_num = <yours>,
you are tuning something already tuned and your comparison needs a different baseline. If it
says will use default config, you have found real headroom.
Two details in that log line worth reading carefully:
padded_M — where this field appears, aiter bucketed M before the lookup, so the M you
tuned is not necessarily the M that hit. Do not generalize that into coverage. Padding is a
property of the entry point, not of aiter: a8w8_blockscale_bpreshuffle matches exactly on
(M,N,K) with no interpolation, so there a row tuned at M=64 serves M=64 and nothing else.
Earlier versions of this document claimed a sparse grid is safe because "a row tuned at one M
serves a range of nearby M" — true for the generic bf16_tuned_gemm path shown above, false for
the blockscale paths, and it cost a run. See the reproduced case below.
libtype — which backend won. This tells you whether your Triton work is even in play
for that shape, or whether hipBLASLt is beating it.
The shape list is the tune — harvest it, never guess it
The same log that reports engagement also tells you, for free, exactly which shapes the workload
dispatches: every miss names one. Take that reading on the untuned baseline you are already
running, and let it write your tuner input.
AITER_LOG_TUNED_CONFIG=1 <your serve command> 2>&1 | tee /tmp/aiter_dispatch.log
# drive the REAL benchmark workload, not a smoke test — the M distribution is the thing you want
grep -o 'shape is M:[0-9]*, N:[0-9]*, K:[0-9]*' /tmp/aiter_dispatch.log | sort -u
The served M values are chosen by the engine, not by you: decode M follows the CUDA-graph capture
buckets and prefill M follows the chunk size. A grid that looks sensible on paper — a few powers of
two, one decode point, one prefill point — can miss every single one of them.
Reproduced (gfx950, sglang + aiter, fp8 a8w8 blockscale bpreshuffle, 85.1% of GPU time). A
per-shape tune was run over M ∈ {1, 64, 1024, 16384} across the four served (N,K) families. The
engine dispatched decode M ∈ {2..512} and prefill chunks M ∈ {~7168..15360} — no intersection.
Result: 0 is tuned on cu_num hits against 272 not found tuned config. The tuned artifact
was never the code that ran. The tuner had genuinely found a 1.28x and it was unreachable as
tuned; the failure was in binding, not in search. The e2e A/B then measured +1.18%, which was box
drift, and anyone not counting hits would have banked it as a win.
The shape list is the engine's property, not the model's. Same model, same workload, same node,
two sglang versions: 0.5.12 dispatched 348 distinct shapes over 5(N,K) families; 0.5.15
dispatched 444 over 12, including prefill M values (7171, 10260, 11945, 13096, 15128, 15176,
15284, 15380, 16276) absent from the first list entirely.
The cause is visible in the engine's own knobs: 0.5.12 has a single cuda_graph_bs, while 0.5.15
splits it into cuda_graph_bs_decode and cuda_graph_bs_prefill, so prefill is now captured at its
own batch sizes. Together with chunked_prefill_size, those knobs — not the model, and not aiter —
are what fixes the M distribution you must cover. Carrying a harvested list across an engine upgrade
silently reintroduces the miss it was built to prevent, and changing either knob invalidates it too.
Re-harvest.
So treat the shape list as a derived artifact, and gate on binding before spending real time:
Harvest the dispatched (M,N,K) set from the baseline log, as above.
Tune those rows.
Re-run the server briefly and require is tuned on cu_num > 0, with misses going to 0 on the
shapes you tuned. Two minutes here rejects an unreachable tune before an e2e A/B spends an
hour measuring nothing.
Only then measure.
Step 3 is not bookkeeping. A tune that fails to bind returns a plausible number, because the
drift it measures is the same size as the win you were hoping for.
2b. Engagement is not selection — the hit can bind to the wrong kernel
A hit count proves the row was found. It does not prove the kernel in that row ran. On some
aiter builds the serving wrapper reads only libtype from the tuned row and drops kernelName,
so a row that says "cktile instance #11" deploys as "cktile, whatever cktile picks by default".
The log still prints is tuned on cu_num. Every gate in §2 passes. And the library's internal
default can be far worse than the default you displaced.
Measured (gfx950, a8w8_blockscale_bpreshuffle, M=15104 N=34816 K=5120, 43.8% of the dominant
head), on an image whose wrapper drops kernelName:
path
time
TFLOPS
production wrapper, no tuned config — what the server did before
4165.0 µs
1293
production wrapper, tuned row installed (libtype=cktile)
7733.7 µs
696
tune entry point, cktile kernelId=11 — what the tuner measured
2646.5 µs
2035
The tuner was right: that kernel is 36% faster than the default, reproducibly, and its recorded
time matches to 0.2%. Deploying it made the op 85.7% slower than doing nothing, because
selecting a library without an instance is not selecting the instance you tuned. Applied across the
shapes carrying ~90% of the head: 290 engagement hits, and e2e throughput −6.48%.
Then the same CSV, same node, same tuner, on a build whose wrapper passes kernelName:
production wrapper 2655.8 µs / 2028 TFLOPS, matching the tuner's recorded 2655.3 µs to 0.02%,
and e2e +23.88% (TTFT −30.8%, TPOT −18.0%). Nothing about the tuning changed. The only
difference was whether the runtime could honour a kernel selection.
So add one gate between "it binds" and "measure e2e" — does the production path reproduce the
tuned number?
# Time the PRODUCTION entry point, the one the server calls, against the tuner's recorded `us`
# for that same row. Not the *_tune entry point, and not --run_config.
op.gemm_a8w8_blockscale_bpreshuffle(x, w_shuffled, x_scale, w_scale)
If it does not land within noise of the CSV's us, the selection is not binding, whatever the hit
count says. This is the whole gate: one shape, two numbers, before you spend an hour on an A/B.
Do not substitute the tuner's own --run_config for this. Measured on the broken build,
--run_config reported 2660 µs — the tuned number, not the 7733 µs the server gets — because it
does not go through the production wrapper. A vendor verification tool that agrees with the vendor
tuner tells you nothing about the vendor runtime.
The static form of the same check, useful before you tune at all:
A=/sgl-workspace/aiter/aiter/ops/gemm_op_a8w8.py
# 1. does the production op even accept a kernel name? (0 = it cannot be told what to run)
grep -A12 'def gemm_a8w8_blockscale_bpreshuffle_cktile(' "$A" | grep -c kernelName
# 2. read the wrapper's branches yourself -- see the trap below before automating this
sed -n '/^def gemm_a8w8_blockscale_bpreshuffle(/,/^def [a-z]/p' "$A" | grep -nE 'libtype|kernelName'
Trap: "does kernelName appear in the wrapper" is not the test, and it returns a false OK on the
broken build. The broken wrapper does contain kernelName=kernelName — in its asm branch —
while the ck and cktile branches call with positional args only:
if libtype == "cktile":
return gemm_a8w8_blockscale_bpreshuffle_cktile(XQ, WQ, x_scale, w_scale, Y) # no kernel
elif libtype == "ck":
return gemm_a8w8_blockscale_bpreshuffle_ck(XQ, WQ, x_scale, w_scale, Y) # no kernel
elif libtype == "asm":
kernelName = config["kernelName"] # only here
Test the signature of the op your row names, and the forwarding in that specific branch.
env-setup/audit_tools.sh does this and is validated against both images — BROKEN on
v0.5.12/a6bb49937, OK on v0.5.15.post1/9127c94a1. Run it before you tune:
BROKEN sig-cktile gemm_a8w8_blockscale_bpreshuffle_cktile takes NO kernelName …
ACTION update-aiter predates 7136b240e (#3075); need >= v0.1.15 …
Read the source, not inspect.signature: these ops are JIT-wrapped and report (*args, **kwargs)
until the module is built. After the first call aiter prints its real signature on a
type hints mismatch, override to --> line, which is the other place to look.
First release containing it: v0.1.15 (2026-05-29). v0.1.14 (2026-05-15) does not. So the
floor is aiter ≥ v0.1.15, and below it a per-shape blockscale GEMM tune is not deployable.
This gate is aiter's alone — the engine version is irrelevant to it. Grepping the whole sglang
tree for kernelName, libtype and AITER_CONFIG_GEMM returns 0 matches, in both v0.5.12 and
v0.5.15.post1. sglang calls gemm_a8w8_blockscale_bpreshuffle(XQ, WQ, x_scale, w_scale) from
srt/layers/quantization/fp8_utils.py and takes no part in choosing a kernel; the wrapper that
drops kernelName lives in aiter/ops/gemm_op_a8w8.py. So do not reason about deployability from
an image tag or an engine version. Ask aiter.
Builds measured (the image tags are only where these aiter commits were found):
aiter commit
git describe
reported version
has 7136b240e
found in
a6bb49937 (2026-04-29)
v0.1.12.post1-150-g…
0.1.12.post2.dev150+…
NO — 169 commits short
sglang v0.5.12-rocm720-mi35x, GEAK CI default
9127c94a1 (2026-06-25)
v0.1.16-110-g…
0.1.17.dev110+…
yes
sglang v0.5.15.post1-rocm720-mi35x
0ba802e2d (2026-07-27)
v0.1.18-46-g…
—
yes
clone_libs.sh pin
The two are bundled in practice, which is why they are easy to conflate, but they vary
independently — and the tag is not even a reliable proxy for the aiter date: the v0.5.12 image's
aiter reports …+ga6bb49937.d20260516, i.e. built 2026-05-16 from a 2026-04-29 commit.
Check it, cheapest first:
# 1. version screen. Anything whose release component is below 0.1.15 is certainly broken.
python3 -c "import aiter._version as v; print(v.__version__)" # e.g. 0.1.17.dev110+g9127c94a1
# 2. authoritative, and sglang images ship the git tree:
git -C /sgl-workspace/aiter merge-base --is-ancestor 7136b240e HEAD \
&& echo "kernelName dispatch PRESENT" \
|| echo "BROKEN: aiter predates #3075 -- update to >= v0.1.15 before tuning blockscale GEMM"
Two traps in reading the version string. It is a PEP 440 dev version — 0.1.17.dev110+g9127c94a1
means 110 commits past v0.1.16, heading for 0.1.17, so it is newer than v0.1.16, not older
than v0.1.17. And there is a genuinely ambiguous band: commits between the fix (2026-05-21) and the
v0.1.15 tag (2026-05-29) contain it while still describing as v0.1.14-rc0-NN. In that window the
version cannot answer the question — use check 2, or the behavioural check above.
The floor is necessary, not sufficient. The commit title says blockscale, and it means it. Even
at v0.1.16+, only 19 production ops accept kernelName: the a8w8 blockscale family
(_ck, _cktile, _bpreshuffle_ck, _bpreshuffle_cktile, _bpreshuffle_asm), a4w4 blockscale,
the a8w8/a16w16 asm entry points, and the MoE stage ops. Roughly 53 others do not — including the
plain, non-blockscale a8w8 and a8w8_bpreshuffle paths. Passing the version check does not license
skipping the per-op check.
On a build below the floor, asm is the only family whose tuning is deployable. A tuned table that
is mostly ck/cktile is then not a small loss but an active regression, and the right move is
to raise the runtime, not to keep tuning. If you cannot change the image, restrict the deployed table
to libtype=asm rows and expect a correspondingly small ceiling.
One side effect of the same commit: it also added strict tuned-CSV validation, so at ≥ v0.1.15 a
malformed or schema-mismatched tuned CSV tends to be rejected outright rather than silently ignored.
That is an improvement, and it softens the §3 warning below — on these builds a tuner/runtime
generation mismatch is more likely to raise than to become a silent miss.
Corollary for reading a no-win. "The tuner found nothing, so the default is already good" and
"the tuner found something the runtime cannot execute" produce the same e2e number and opposite
conclusions. The first says stop; the second says fix the runtime and a large win is waiting. Only
the production-path timing above distinguishes them, and the original 27B run — which concluded
"default CK heuristic already near-optimal in the live serving context" — had it backwards by 30
percentage points.
3. Get the tuner (vllm image only)
The images differ in a way that changes the workflow:
sglang ships aiter as a git checkout at /sgl-workspace/aiter — tuners present, tune
in place.
vllm ships aiter as a wheel — configs/*.csv present, but no gradlib/, no csrc/.
Wheels carry the runtime, not the build tree.
For vllm, supply source matching the installed wheel:
Resolve the tag from the installed version rather than copying a hash. Do notpip install the checkout over the wheel — run the tuner from the source tree and let it
write into the installed package. A tuner from a different aiter generation can emit columns
the installed runtime does not read, which fails as a silent miss rather than an error.
4. Run the tuner
First, confirm which tuner your op uses. This section documents gradlib's
gemm_tuner.py, which serves dense bf16 only. Every quantized GEMM goes through a
different per-op tuner with a different entry point, different flags, a different schema, and
a different cost — none of which is what you read below. Getting this wrong is not a small
detour: you will run a tuner that never touches the op you care about.
your op
tuner
this section applies?
dense bf16
gradlib/gradlib/gemm_tuner.py
yes
fp8 / int8 / fp4 / batched / MoE
csrc/ck_gemm_*/..._tune.py, or aiter/utility/pretune.py
no — see ../tuning-ck/ §3 and §4b below
4b. If it is a per-op quantized tuner, these are the differences that bite
Verified on aiter d9e5ef7c / gfx950 against
csrc/ck_gemm_a8w8_bpreshuffle/gemm_a8w8_bpreshuffle_tune.py, the FP8 pre-shuffled path an
FP8 SGLang deploy actually dispatches:
gradlib (below)
per-op quantized tuner
flags
--input_file / --tuned_file
-i / -o
--libtype
no such flag
exists, default=["all"], choices all, asm, ck, cktile, flydsl
off by default; -k/--splitK is store_true (base_tuner.py:130)
cost, 2 shapes
120 s (gfx950)
≈50 min across 4 libtypes — plan for ~14× the gradlib figure
Three of those change what you do:
There is no torch floor, so build one. The checklist item below tells you to require a
torch row as the baseline the winner must beat. On this tuner that row cannot exist. The
real floor is the CK kernelId=0 fallback — what the op dispatches when your config is
absent — and you have to measure it yourself, on a harness you control. Do not substitute the
tuner's own before/after for it; see Rule 3b/5 in ../tuning-core/measurement.md and the
cold-timing warning in ../tuning-ck/ §3c.
Split-K is off unless you ask. The schema has a splitK column, which makes the search
look fully explored while it holds a zero in every row. With the flag unset the tuner never
races a single split-K candidate, for either the ck or cktile path
(useSplitK = args.splitK, gemm_a8w8_bpreshuffle_tune.py:631). Pass -k for
tall-skinny-M or long-K shapes; budget roughly 4× the runtime.
Scope the sweep from the right number. Anyone planning a 23-shape run from the 120 s /
214 s figures below will discover a multi-hour job. Time two shapes first, then extrapolate.
4a. gradlib's gemm_tuner.py (dense bf16)
Input is a CSV of shapes with the same key columns minus the results:
Do not pass --indtype, and do not look for --libtype. Both appear in every write-up
of this tuner, including earlier versions of this document, and neither works at the aiter
commits currently in play (0ba802e2, 2026-07-27, which clone_libs.sh pins; and
a6bb4993, 2026-04-29, which the sglang image ships). Verified on both.
--indtype raises. gemm_tuner.py:121 rewrites the argument in place from the CLI string
to a torch dtype, and GemmTuner.pre_process then looks the string up again:
The two files disagree about the type of one attribute, so this is aiter's bug rather than
yours — but there is no invocation with --indtype that runs. Put the dtype in the CSV's
dtype column instead (the schema above already has one) and omit the flag: args.indtype
stays None, the conversion is skipped, and the column is used.
--libtype does not exist as a flag at all. libtype is a column in the tuned CSV and a
filter argument to chip_info.build_tune_dict, not a command-line option. The cross-backend
race still happens and the winning backend is still reported per row — you just cannot
restrict it from the command line. So this remains the reason aiter is the integration point
for the per-language skills rather than a competitor to them, and the table below is how to
read the libtype column of the output, not a menu:
libtype
what won
authoring-time tuning lives in
hipblaslt
a hipBLASLt solution index
../tuning-hipblaslt/
triton
one of aiter's Triton kernels
../tuning-triton/
flydsl
a FlyDSL kernel
../tuning-flydsl/
ck / cktile
a Composable Kernel instance
../tuning-ck/
asm
a hand-written assembly kernel
—
skinny
a narrow-GEMM specialization
—
torch
the torch fallback — your floor
—
A torch row is the baseline the winner had to beat, measured by the same harness, and it is
what aiter falls back to on a miss. This holds for gradlib only — the quantized per-op
tuners have no torch candidate, and their floor is the CK kernelId=0 fallback, which you must
measure yourself (§4b).
Measured on gfx950, 2 shapes, all backends: 120.2 s (gfx942, 2 shapes: 214 s). Not
transferable to the per-op quantized tuners, which ran ~50 min for the same two shapes. Cost
scales with the shape count, so scope the shape list — but scope it by harvesting what the
workload actually dispatches (§2), never by thinning a generated sweep. Capturing from the live
workload is a correctness requirement, not a cost optimization: a shape you did not harvest is a
shape that will fall back. Per-engine capture details are in ../tuning-in-vllm/ and
../tuning-in-sglang/.
Useful flags: --compare benchmarks before and after and prints the delta;
--update_improved --min_improvement_pct N only writes rows that improved by at least N%,
which is the built-in defence against enshrining a within-noise "win".
That defence has a hole, and it is not theoretical. Shapes not already present in the tuned
CSV are classified NEW and written regardless of the measured comparison — including when the
comparison shows a slowdown. Verified on gfx950 for fp8 M=1, N=7168, K=8192: measured 0.82x,
printed OK, improvement reported as N/A, row written. Read the Speedup column yourself.
Worked through in ../tuning-ck/ §3b, which is also where the per-op CK tuners live.
Output carries your keys and a correctness column:
Worth comparing column by column, because which parts moved is the useful part:
gfx942
gfx950
gfx / cu_num — the lookup key
gfx942 / 304
gfx950 / 256
4096³ solution index
198969
440518
4096³ macro-tile
MT256x224x64
MT256x256x64
4096³ throughput
563.9 TFLOPS
1327.5 TFLOPS
M=1 macro-tile
MT16x16x512
MT16x16x512 — unchanged
M=1 bandwidth
2446.7 GB/s
3548.4 GB/s
The solution index is a different number on a different part, and the ISA942/ISA950 field
in the kernel name says why: these are separately compiled kernels, so an index copied across
parts names something unrelated or nothing. The macro-tile, by contrast, transfers as a
family — M=1 wants 16x16x512 on both, and the square shape wants a 256-wide tile on both,
merely 256 deep rather than 224. Read the tile to learn what the problem wants; re-derive the
index per part.
Read the two rows together: the 4096³ winner uses a 256x224x64 macro-tile at 563.9 TFLOPS;
the M=1 winner uses 16x16x512 at 2.45 TFLOPS but 2447 GB/s. Opposite tile shapes,
opposite figures of merit. Judge the decode row on bandwidth, not FLOPS
(../tuning-core/measurement.md). Check err_ratio on every row before deploying — the
tuner records it precisely so you can reject a fast-but-wrong solution.
5. Deploy to the directory the runtime reads — not the one it names
This is the step most likely to silently do nothing.
The log line points at /tmp/aiter_configs/bf16_tuned_gemm.csv, and
AITER_CONFIGS.AITER_CONFIG_GEMM_BF16_FILE returns that same path. Both are true and both
are traps: /tmp/aiter_configs/ is a derived cache, regenerated on import by merging
aiter/configs/*.csv with aiter/configs/model_configs/*.csv.
Measured — appending correct gfx942/304 rows directly to the /tmp file:
rows before import: 0 # (after writing 2 rows, then importing aiter)
rows after import: 0
The edit is erased and the lookup still reports not found tuned config. Writing to the
path the runtime names accomplishes nothing.
Deploy into the installed package's config directory instead. There are two ways, and
prefer the drop-in:
# PREFERRED — a new file in model_configs/, which the import step merges in
C=/usr/local/lib/python3.12/dist-packages/aiter/configs/model_configs
cp /tmp/tuned.csv "$C/a8w8_bpreshuffle_tuned_gemm_mymodel.csv"
This is what upstream itself does — aiter already ships
a8w8_bpreshuffle_tuned_gemm_{glm47,glm5.2,minimax_m3,…}.csv there. It produces a
few-line reviewable diff, exports cleanly as a patch, and A/B is mv the file away. Nothing
shipped is modified.
# WORKS, but mutates a shipped file — use only if the op has no model_configs/ path
P=/usr/local/lib/python3.12/dist-packages/aiter/configs/bf16_tuned_gemm.csv
cp "$P" "$P.bak" # always, so you can prove a clean A/B
tail -n +2 /tmp/tuned.csv >> "$P" # skip the header row
Appending rewrites a 598-row vendored file, so your diff is indistinguishable from a vendor
change, and reverting depends on a .bak that lives outside version control.
Then the same lookup reports:
found padded_M: 4096, N:4096, K:4096 is tuned on cu_num = 304 ... libtype is hipblaslt
That transition is the cheap engagement check. Two corrections to how it is usually stated,
both of which have cost real debugging time:
The transition is only not found tuned config → is tuned on cu_num = 304when
AITER_LOG_TUNED_CONFIG=1 is set. The miss line prints unconditionally
(gemm_op_a8w8.py:456, :519); the hit line sits behind if AITER_LOG_TUNED_CONFIG:
(:449, :509). Without the flag the transition is not found tuned config → silence,
so grep -c "is tuned on cu_num" returns 0 on a deploy that is working perfectly. If you
cannot set the flag — because the environment is the measurement contract — then use the
disappearance of the miss line, which needs no flag, and keep an untuned shape as a negative
control.
And a config hit is not the strongest available proof, despite how this is usually
phrased. It shows a lookup succeeded, not that the tuned kernel executed — and on some
builds those differ, which is exactly the §2b failure where the wrapper reads libtype and
discards kernelName. Kernel identity in a profile is strictly stronger and does not go stale
when the library changes its logging. Use the log line as the fast loop; confirm anything you
ship with a trace (../tuning-core/engagement_verification.md, form 4).
Generalize the rule: find the deploy target empirically. Write a row, import fresh, and
check whether it survived and whether the lookup hits. Do not infer the target from a config
constant or a log path.
The derived cache has a more dangerous sibling: the merge writes back
The trap above costs you a change that does not take effect, which is loud and cheap. The same
merge step can also edit the source file you are about to export as a patch, which is
neither.
If a shape appears in more than one config source — say your drop-in duplicates a row aiter
already ships — the merge does not error and does not prefer yours. It sorts by the us
column and keeps the lowest (aiter/utility/base_tuner.py:300-303), rewrites
aiter/configs/model_configs/*.csv in place, and on the version in the SGLang image raises a
RuntimeError asking you to re-run. Re-running then succeeds against the file it just
modified.
Two things make this the worst failure mode in this document:
The arbiter is the one number you should not trust.us comes from the tuner's own
timing, which is warm-cache and was measured 23% optimistic on one shape in this workload
(../tuning-ck/ §3c). A row that loses on cold-cache reality can win the merge.
It corrupts the artifact, not the run. Every other trap here produces a null result.
This one produces a wrong patch that reproduces — because the file you export already
contains the substitution.
Defend against it: git diff (or md5sum) your config files after the first import that
follows a deploy, not just after your edit. Keep your rows in a uniquely named drop-in file
so duplicates cannot arise, and if a shape genuinely overlaps a shipped row, decide the winner
yourself with a cold measurement and delete the loser rather than letting us arbitrate.
The flag is mandatory, not decorative — without it the hit line is never emitted and this
command returns 0 no matter how well the deploy works (§5).
Expect a non-zero count. Do not expect the count to track how often the shape runs: the
lookup is wrapped in functools.lru_cache(maxsize=1024) keyed on raw M
(gemm_op_a8w8.py:360, :401, :416, :465), so each distinct M logs once and then goes
quiet, with sporadic re-logging as eviction thrashes. Log-line counts measure the diversity
of M, not call frequency. Measured on Qwen3-8B decode: 617 distinct M values, and the single
most valuable shape (M=64, executing ~110k times) logged twice — while a diffuse prefill
band logged hundreds of times for far less wall clock. Ranking targets by miss count inverts
the priority order. Rank by measured time.
Zero means either the flag is unset or the deploy did not take, in that order of likelihood.
Distinguish them with the miss line: if it still names your shape, the deploy really did not
take; if it has gone quiet, the lookup is fine and the logging is what is missing.
Only after engagement is established does an end-to-end timing mean anything. Restore from
.bak and re-measure to confirm the delta is attributable to your rows and not to drift — and
note that any config change requires a process restart, so your noise floor is the
restart-to-restart spread, not the repeat-the-benchmark spread. These differ by ~26× on this
workload; see ../tuning-core/measurement.md Rule 3b.
7. The MX (MXFP4/MXFP8) ops on gfx950: a different table, and four traps
The microscaled ops are the gfx950-only surface and they do not go through
the CSV mechanism the rest of this skill describes. They are Triton kernels
configured from JSON under aiter/ops/triton/configs/, keyed on arch and shape
rather than on cu_num, and read through a per-op _get_config(M, N, K) rather
than through the tuner. §1–§6 do not apply to them.
MXFP8 is not here. On both shipped images the MX surface is FP4-only:
gemm_afp4wfp4, gemm_a16wfp4, gemm_a8wfp4, batched_gemm_afp4wfp4,
fused_moe_mxfp4, dynamic_mxfp4_quant. There is no gemm_afp8wfp8 and no
dynamic_mxfp8_quant. Newer aiter source has them, but mostly bound to
gfx1250 code objects, which do not load on MI355. For an MXFP8 GEMM on this
part, use CK — see ../tuning-ck/SKILL.md §2c. Where FP8 appears on the gfx950
MX path it is as the activation side against FP4 weights (gemm_a8wfp4).
Four faults worth knowing before writing anything against these ops. Each was
hit while building corpus cases, and none of them reports the actual problem:
_get_config does not take the same K in every op. The batched op derives
its shapes from the packed tensors it is handed, so it works in K/2; the
unbatched op takes logical K.
gemm_afp4wfp4._get_config(M, N, K) # logical K
batched_gemm_afp4wfp4._get_config(M, N, K // 2) # packed K
The knob that overruns is SPLITK_BLOCK_SIZE: the op sets it to 2*K and then
strides by SPLITK_BLOCK_SIZE // 2 over a buffer of K/2 packed bytes, so the
logical-K call strides twice the buffer. BLOCK_SIZE_K comes back at 256 either
way on most shapes, which is why this is hard to spot by diffing configs.
Since it is an out-of-bounds read, the symptom is undefined and not stable
enough to debug by observation: the same mistake gives Memory access fault by GPU node-5 ... Reason: Unknown (no traceback, process dies) at K=4096, a silent
nan at K=256, and a plausible in-tolerance answer on a rerun of that same
call. A clean run does not mean the convention was right. Nothing in any of the
three outcomes mentions K.
Ops with an output parameter return None.gemm_a8wfp4,
batched_gemm_afp4wfp4 and fused_moe_mxfp4 all write into a buffer you pass
and all have docstrings promising Returns: torch.Tensor. Read the buffer.
batched_gemm_afp4wfp4(y=None) raises.y is typed
Optional[torch.Tensor] = None; the body does By, _, _ = y.shape before any
check. It is required.
In fused_moe_mxfp4, BLOCK_SIZE_M changes the input. The token-to-block
map from moe_align_block_size_triton is built fromBLOCK_SIZE_M. Sweeping
that knob against a map computed once from the default config feeds candidates a
token layout built for a different block size, and nothing range-checks it.
Rebuild the alignment per config.
The payoff for going through all that is that these tables are thin —
gfx950-MOE-MX_FP4.json is 689 bytes and holds one config for every shape — so
there is a lot on the table. The corpus cases find 34.8% on gemm_a8wfp4 and
73.1% on gemm_mxfp4_batched, the largest uplift in the corpus.
Checklist
get_gfx() / get_cu_num() recorded; shipped rows checked for your key
aiter ≥ v0.1.15 for blockscale GEMM tuning — i.e. contains 7136b240e (§2b). Below that
floor, stop and raise the runtime; only asm rows are deployable
runtime can select a kernel, not just a library — your production op takes kernelName
and the wrapper passes it (§2b). The version floor covers blockscale only; 53 ops still do not
AITER_LOG_TUNED_CONFIG=1 baseline taken before any change — and if the environment is
frozen so you cannot set it, the flag-free substitute is in place instead: miss-line
disappearance plus an untuned negative control (§5)
shape list harvested from that baseline log — not authored, not a generated grid, and
re-harvested for this build, not carried over from another image (§2). Targets ranked
by measured time, never by log-line count, which measures diversity of M (§6)
correct tuner identified for the op — gradlib for dense bf16, a per-op
csrc/ck_gemm_*/..._tune.py for anything quantized (§4)
tuner source version matches the installed aiter.__version__
floor established: a torch row for gradlib, or — for per-op quantized tuners, which have
no torch candidate — the CK kernelId=0 fallback measured on your own harness (§4b)
-k/--splitK decision made deliberately; it is off by default and the zero-filled
splitK column does not mean the space was searched (§4b)
err_ratio / errRatio inspected per row before deploying
every winner re-timed cold on a harness you control before it ships — the tuner's us
is warm-cache and was 23% optimistic on one shape here ( §3c)
../tuning-ck/
rows deployed as a uniquely named drop-in under the packageconfigs/model_configs/,
not appended to a shipped file and not written to /tmp/aiter_configs/ (§5)
config files verified unmodified after the first post-deploy import — the merge
silently resolves duplicate shapes by us and rewrites your source file (§5)
.bak kept; lookup transition observed
engagement pre-gate passed before any e2e A/B — is tuned on cu_num > 0 with
AITER_LOG_TUNED_CONFIG=1 set, misses → 0, and kernel identity confirmed in a trace for
anything you intend to publish
selection-fidelity pre-gate passed — the production entry point reproduces the tuned
row's recorded us (§2b). Not *_tune(kernelId=…), not --run_config
incumbent floor measured through the production path, so a "null" can be told apart from an
unexecutable win (§2b corollary)
A/B run with the reference leg repeated (ref → cand → ref) so drift is measured rather
than assumed. A config or CSV change cannot take effect without a process restart, so
interleave across restarts and compare distributions — the same-session spread
understates the real floor by ~26× here (../tuning-core/measurement.md Rule 3b)
correctness gated against a same-config control (ref vs ref), requests issued one at a
time — see ../tuning-core/correctness_gates.md
