Repository navigation
[Config] Extend the DeepSeek-V4 a8w8 blockscale GEMM tunings for gfx950 - #5485
Conversation
github-actions
Bot
changed the title
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
PR title tags & labels: |
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
PR title tags & labels: |
|
Reopened. Disregard my close note above: #5473 turned out to be permanently un-reopenable, so this PR is the one to review. GitHub rejects reopening #5473 under two mutually exclusive rules. Restoring its recorded head gives Content here is unchanged from what I described: 734 rows added and 4 updated in |
|
Closing in favour of a cleaner branch: the shape additions and the four retuned rows are now separate commits, so the new rows can be reviewed apart from the changes to existing ones. Same content, same measurements. |
1f780d8 to
8c3921d
Compare
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
PR title tags & labels: |
|
@valarLip this carries on from your review on #5473 ( Your point is applied, and chasing it down showed I had a second problem: Split into two commits so the new shapes can be read apart from the four rows that touch existing entries. Those four were re-timed on hardware rather than compared across tuning runs; 15 other differing picks turned out to be tuner noise and are left alone. |
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
PR title tags & labels: |
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, and 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, each falling back to a kernel chosen without knowledge of the tile. Per-shape, 40 shapes sampled across the M range of all six families, measured against the current table where all 40 fall back: median 2.79x, mean 2.94x, max 5.86x, 50860us -> 11700us in aggregate; the gain grows with M. One shape, M=64 N=7168 K=3072, is 0.94x, within run-to-run spread at 12us. In a captured prefill trace the blockscale-family GEMM time per step per rank falls from 500.1ms to 244.8ms, with the untuned ck::kernel_gemm_xdl_cshuffle_v3 fallback going from 171 calls at 2059us to 46 at 555us. End to end on DeepSeek-V4-Pro, MI355X x8, TP1/DP8, FP8 KV, MTP-3, concurrency 64, full 3600-second window with everything else byte-identical: 27,003,552 -> 30,319,229 tokens/$ (+12.3%) and 7.055 -> 8.654 P90 interactivity (+22.7%).
At M = 4, 128, 256 and 512 the best pick moves off the ck kernels. Timed here
against the rows they replace, both kernels back to back in one process on the
same tensors with the arms interleaved, medians over 7 repeats of 50 iterations
after 20 warmup calls, on MI355X gfx950 cu_num=256:
M current proposed speedup
512 78.01us 61.75us 1.26x
128 25.17us 21.11us 1.19x
4 17.15us 16.22us 1.06x
256 42.07us 40.50us 1.04x
Tuning also produced a different pick for 15 further shapes already in the file.
Re-timing them the same way put every one within 1% of the row already there, so
they are left untouched: 13 resolve to the identical ck_tile tile config
192x256x128_4x2x1_16x16x128_intrawave_0x1x0 and differ only in a trailing
variant index, meaning the spread between the two tuning runs was noise rather
than a kernel difference. No shape regressed.
8c3921d to
10aada0
Compare
…ing fixes, fp8 MQA logits split-k and DSv4 tunings (#5573, #5295, #5558, #5603, #5627, #5485) (#5638) Cherry-picks six already-merged `main` PRs onto `release/v0.1.22` for the `v0.1.22.post1` post release. | PR | `main` commit | Backport commit | What | |---|---|---|---| | #5485 | `9252f4672` | `155534984` | Extend the DeepSeek-V4 a8w8 blockscale GEMM tunings for gfx950 (tuning CSV only) | | #5573 | `22d2c7c91` | `6b23ba866` | Pad the MXFP4 A4W4 MoE sort extent to a block_size multiple (fixes a HIP illegal memory access) | | #5295 | `972c8e1fd` | `7d68b0edb` | Skip invalid expert IDs in MoE sorting | | #5558 | `3fdfca11e` | `dd83a9d17` | Fix MoE routing kernel compile failure | | #5603 | `a84bd368c` | `a96461997` | Add split-k support for fp8 MQA logits on gfx950 | | #5627 | `f5ed7dc54` | `a41214712` | Follow-up to #5603: drop chunking when summation folding is unavailable (fixes Triton 3.6 compile) | Original PRs: - #5485: #5485 - #5573: #5573 - #5295: #5295 - #5558: #5558 - #5603: #5603 - #5627: #5627 To be published as `v0.1.22.post1` once merged (tag on the merge commit, release automation builds the wheel set).
2 participants
Summary
Two commits against
aiter/configs/model_configs/dsv4_a8w8_blockscale_tuned_gemm.csv: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_filemergesmodel_configs/withconfigs/at load time, so per-model rows belong here.This continues #5473, which I closed by accident with a force-push of a commit
that had lost its parent. GitHub refuses to reopen it: the reopen check runs
against that PR's recorded head sha, which is the parentless commit, so no state
of the branch can satisfy it. The review there still applies and is answered
here.
Commit 1: the missing shapes
The file already covers four of the six
(N, K)families this model drives, butonly at a coarse M ladder, and carries no entry at all for the two vocabulary
projections:
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 defaultERR_RATIOon gfx950 withcu_num=256, matching the rest of the file. M valuesfollow 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 configfor everydistinct
(M, N, K)DeepSeek-V4 issues; in one 3600-second AgentX run on MI355Xthat 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
usvalues that come from two separate tuning runs, I re-timed all 19on 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 theckkernelsto the
ck_tileaqrmvariants: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_tiletile config192x256x128_4x2x1_16x16x128_intrawave_0x1x0and differ only in a trailingvariant 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 thesame image: once against the current table, where all 40 sampled shapes report
not found tuned config, and once with these rows present. 40 shapes sampledacross the M range of all six families.
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=3072at 0.94x. It is the only one of the40 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:
ck::kernel_gemm_xdl_cshuffle_v3352.2 ms, n=171, 2059 us/callGEMM time per prefill step halves, and it is the untuned
kernel_gemm_xdl_cshuffle_v3fallback that disappears as the work moves ontotuned 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:
+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=256only. The shipped table is gfx950/256throughout so this matches it, but these rows say nothing about other
architectures.