[Config] Extend the DeepSeek-V4 a8w8 blockscale GEMM tunings for gfx950

[LiuYinfeng01]

Summary

Two commits against aiter/configs/model_configs/dsv4_a8w8_blockscale_tuned_gemm.csv:

1. Add 734 rows for shapes the file does not cover.
2. Retune 4 rows that it does, all N=65536 K=1536.

Addressing @valarLip's review on #5473 ("move new configs to model_configs"): my
first attempt put these in the global aiter/configs/a8w8_blockscale_tuned_gemm.csv.
That was wrong twice over. The per-model file already exists, and 81 of the rows
I proposed collided with keys already in it. get_config_file merges
model_configs/ with configs/ at load time, so per-model rows belong here.

Supersedes #5473 and #5485. I closed #5473 by accident with a force-push of a
commit that had lost its parent; GitHub read it as unrelated history and will
not reopen it. #5485 carried both changes in one commit.

Commit 1: the missing shapes

The file already covers four of the six (N, K) families this model drives, but
only at a coarse M ladder, and carries no entry at all for the two vocabulary
projections:

N	K	before	added	after
7168	7168	61	594	655
6144	7168	15	39	54
7168	3072	15	39	54
65536	1536	15	39	54
129280	7168	0	16	16
16160	7168	0	7	7

Pure addition: the 310 existing rows are byte-identical, no key is duplicated,
every added row has errRatio = 0.0. Tuned with aiter's own tuner at the default
ERR_RATIO on gfx950 with cu_num=256, matching the rest of the file. M values
follow the ladder the file already uses, restricted to the range each family
reaches at runtime.

Without these rows the entry point reports not found tuned config for every
distinct (M, N, K) DeepSeek-V4 issues; in one 3600-second AgentX run on MI355X
that fires 14066 times.

Commit 2: the four retuned rows

Of the 81 colliding keys, tuning produced a different kernel pick for 19. Rather
than trust us values that come from two separate tuning runs, I re-timed all 19
on one MI355X: both kernels back to back in a single process on the same tensors,
arms interleaved, medians over 7 repeats of 50 iterations after 20 warmup calls.

Only 4 are real, all N=65536 K=1536, where the pick moves from the ck kernels
to the ck_tile aqrm variants:

M	N	K	current	proposed	speedup
512	65536	1536	78.01us	61.75us	1.26x
128	65536	1536	25.17us	21.11us	1.19x
4	65536	1536	17.15us	16.22us	1.06x
256	65536	1536	42.07us	40.50us	1.04x

The other 15 landed within 1% of the row already in the file and are left
untouched. 13 of them resolve to the identical ck_tile tile config
192x256x128_4x2x1_16x16x128_intrawave_0x1x0 and differ only in a trailing
variant index, so the spread in the CSV was tuner run-to-run noise, not a kernel
difference. No shape regressed.

These four are worth roughly 22us per occurrence set, far below what an
end-to-end run resolves, so nothing below is attributed to them.

Per-shape effect

Same entry point the model calls, aiter.gemm_a8w8_blockscale, run twice in the
same image: once against the current table, where all 40 sampled shapes report
not found tuned config, and once with these rows present. 40 shapes sampled
across the M range of all six families.

M	N	K	fallback	tuned	speedup
4	6144	7168	25.31 us	21.72 us	1.17x
112	6144	7168	43.64	30.72	1.42x
448	6144	7168	114.00	49.88	2.29x
16384	6144	7168	3458.92	647.58	5.34x
64	7168	3072	12.42	13.15	0.94x
448	7168	3072	58.43	26.25	2.23x
16384	7168	3072	1706.08	362.31	4.71x
7332	7168	7168	1822.48	362.20	5.03x
22832	7168	7168	5620.46	958.78	5.86x
32768	7168	7168	8046.69	1509.39	5.33x
272	16160	7168	225.69	62.01	3.64x
368	16160	7168	295.69	59.81	4.94x
112	65536	1536	59.67	22.06	2.71x
16384	65536	1536	7715.77	2649.00	2.91x
8	129280	7168	160.42	157.02	1.02x
128	129280	7168	609.75	196.51	3.10x

Over the 40 sampled shapes: median 2.79x, mean 2.94x, max 5.86x, and
50,860 us to 11,700 us in aggregate. The gain grows with M, which is the
expected shape of the problem.

One shape regresses, M=64 N=7168 K=3072 at 0.94x. It is the only one of the
40 and the absolute difference is 0.7 us, within run-to-run spread, but it is
listed rather than omitted.

In situ

The same effect on a live run, summing every blockscale-family GEMM kernel per
prefill step per rank from a captured trace:

	GEMM per step	largest contributor
current table	500.1 ms	ck::kernel_gemm_xdl_cshuffle_v3 352.2 ms, n=171, 2059 us/call
+ these rows	244.8 ms	same kernel: 25.6 ms, n=46, 555 us/call

GEMM time per prefill step halves, and it is the untuned
kernel_gemm_xdl_cshuffle_v3 fallback that disappears as the work moves onto
tuned cktile kernels.

End to end

DeepSeek-V4-Pro, MI355X x8, TP1/DP8, FP8 KV, MTP-3, concurrency 64, full
3600-second window, everything else held byte-identical:

Configuration	tokens/$	P90 interactivity	requests	a8w8 fallbacks
current table	27,003,552	7.055	2836	14066
+ these rows	30,319,229	8.654	3140	7180

+12.3% tokens/$ and +22.7% P90 interactivity from a data file alone.

The remaining 7180 fallbacks are M values between ladder points. The lookup
still resolves those, they are simply not exact hits.

Relation to #4664

#4664 added DeepSeek-V4 coverage but left these six families out. The rows here
are diffed against current main, not against that PR.

Scope

Tuned on gfx950 with cu_num=256 only. The shipped table is gfx950/256
throughout so this matches it, but these rows say nothing about other
architectures.

[github-actions]

🏷️ CI Guide

Runs automatically on every PR:

• ✅ Pre-checks (submodule verification, code formatting)
• ✅ Aiter op tests (gfx942 + gfx950)
• ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label	Tests
ci:gfx1250-ffm-triton	Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x	Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
multigpu	Aiter multi-GPU tests on the 8-GPU runner
ci:sglang	SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom	ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full	ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm	vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all	All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 5486 --add-label <label>

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to opt this PR out.

[LiuYinfeng01]

Closing: I should not have opened this. The work continues in #5485, which now carries the same two commits.