tuning-aiter

name	tuning-aiter
description	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:

gfx,cu_num,M,N,K,bias,dtype,outdtype,scaleAB,bpreshuffle,libtype,solidx,splitK,us,kernelName,err_ratio,tflops,bw

The first two columns are the whole story of why tuned artifacts do not travel. A lookup matches on gfx and cu_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:

1. Harvest the dispatched (M,N,K) set from the baseline log, as above.
2. Tune those rows.
3. 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.
4. 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.

The version boundary — know it by name

The fix is a single identifiable commit:

7136b240e102e3b54bc8e960abf59e59f953f5c8   2026-05-21
blockscale gemm: dispatch by kernelName, strict tuned-CSV validation (#3075)

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:

python3 -c "import aiter._version as v; print(v.__version__)"     # -> 0.1.13
git clone https://github.com/ROCm/aiter.git /work_aiter
git -C /work_aiter checkout $(git -C /work_aiter rev-list -n1 v0.1.13)

Resolve the tag from the installed version rather than copying a hash. Do not pip 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
torch in the race	yes, and it is your floor	not selectable — there is no torch candidate
result schema	…,libtype,solidx,splitK,us,kernelName,err_ratio,tflops,bw	…,kernelId,splitK,us,kernelName,tflops,bw,errRatio (base_tuner.py:1546)
split-K	—	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:

M,N,K,bias,dtype,outdtype,scaleAB,bpreshuffle
4096,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False
1,4096,4096,False,torch.bfloat16,torch.bfloat16,False,False

export HIP_VISIBLE_DEVICES=4          # pin an idle GPU
python3 gradlib/gradlib/gemm_tuner.py \
    --input_file  /tmp/untuned.csv \
    --tuned_file  /tmp/tuned.csv

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:

args.indtype = get_dtype(args.indtype)              # "bf16" -> torch.bfloat16
...
self.untunedf["dtype"] = f"dtypes.{_cli_to_dtypes[args.indtype]}"
# KeyError: torch.bfloat16

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:

gfx942,304,4096,4096,4096,False,torch.bfloat16,...,hipblaslt,198969,0,243.7286,Cijk_...MT256x224x64...ISA942...,0.0,563.9,413.01
gfx942,304,1,4096,4096,False,torch.bfloat16,...,hipblaslt,198198,0,13.7208,Cijk_...MT16x16x512...ISA942...,0.0,2.45,2446.71

The same two shapes on gfx950, from the run above:

gfx950,256,4096,4096,4096,False,torch.bfloat16,...,hipblaslt,440518,0,103.5316,Cijk_...MT256x256x64...ISA950...,0.0,1327.51,972.3
gfx950,256,1,4096,4096,False,torch.bfloat16,...,hipblaslt,439841,0,  9.4608,Cijk_...MT16x16x512...ISA950...,0.0,   3.55,3548.41

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 = 304 when 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:

1. 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.
2. 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.

6. Verify, then re-measure

AITER_LOG_TUNED_CONFIG=1 <your workload> 2>&1 | grep -c "is tuned on cu_num"

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 from BLOCK_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)