Repository navigation
Conversation
github-actions
Bot
added
documentation
kevin-mii
force-pushed
the
dsv41-amd-main
branch
from
September 17, 2026 03:31
f5cd9c7 to
7791f6c
Compare
yangyuhuiling
pushed a commit
to ROCm/mori
that referenced
this pull request
Sep 20, 2026
The SGLang-side numbers were taken against a synthetic snapshot of the container image, because the V4.1 code existed on no branch at the time. It does now -- sgl-project/sglang#39857, `kevin-mii:dsv41-amd-main` -- so all three servers were re-run on top of it. The fused wo_b reproduces and is slightly better: fp8 wire +5.6% to +6.0% at bs 4-16 against +5.1% to +5.7% before, bf16 +1.9% to +2.8% against +2.3% to +2.6%. GSM8K 0.917 / 0.912 / 0.918. Added, because it had never been in this section: the standalone GEMM on its own, +2.2% to +2.5% with nothing fused, GSM8K 0.920 / 0.928. And a note that the two results do not add -- `fused_wo_b` runs first and the linear only sees what it declines, so they split `wo_b` rather than stacking on it. Both together is a fourth configuration nobody has measured. **One claim here was wrong and is corrected.** It called bs=1 a built-in control on the grounds that M=4096 is below both floors. The floors are 8192 for bf16 and 2048 for fp8, and the shape log confirms the consequence: the bf16 wire's smallest served M is 7734, the fp8 wire's is 1792. bs=1 is a control for the bf16 column only, and the fp8 column's +3.4% there is a real win that was being dismissed as drift. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
yangyuhuiling
pushed a commit
to ROCm/mori
that referenced
this pull request
Sep 20, 2026
The SGLang-side numbers were taken against a synthetic snapshot of the container image, because the V4.1 code existed on no branch at the time. It does now -- sgl-project/sglang#39857, `kevin-mii:dsv41-amd-main` -- so all three servers were re-run on top of it. The fused wo_b reproduces and is slightly better: fp8 wire +5.6% to +6.0% at bs 4-16 against +5.1% to +5.7% before, bf16 +1.9% to +2.8% against +2.3% to +2.6%. GSM8K 0.917 / 0.912 / 0.918. Added, because it had never been in this section: the standalone GEMM on its own, +2.2% to +2.5% with nothing fused, GSM8K 0.920 / 0.928. And a note that the two results do not add -- `fused_wo_b` runs first and the linear only sees what it declines, so they split `wo_b` rather than stacking on it. Both together is a fourth configuration nobody has measured. **One claim here was wrong and is corrected.** It called bs=1 a built-in control on the grounds that M=4096 is below both floors. The floors are 8192 for bf16 and 2048 for fp8, and the shape log confirms the consequence: the bf16 wire's smallest served M is 7734, the fp8 wire's is 1792. bs=1 is a control for the bf16 column only, and the fp8 column's +3.4% there is a real win that was being dismissed as drift. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
yangyuhuiling
added a commit
to ROCm/mori
that referenced
this pull request
Sep 21, 2026
…ctive (#690) * gemm_ar: MXFP8 (32-wide ue8m0), the quantisation V4.1-Flash actually ships `GemmAllReduceOp` served 1x128 / 128x128 fp8 block scales. DeepSeek-V4.1-Flash quantises 32-wide ue8m0 on both operands (`weight_block_size [32, 32]`, `scale_fmt "ue8m0"`), which is a different operand contract and a different kernel, not a parameter of the old one. CDNA4's `v_mfma_scale_f32_16x16x128_f8f6f4` takes the ue8m0 scales as *instruction operands*, so this form needs no dequantisation arithmetic at all -- no second accumulator, no running rescale. `_BlockScaleK`'s promote/rescale chain exists precisely because a 128-wide block cannot be expressed that way. Two layout results came out of making it fast, both about addresses rather than bytes, and both carry the counters that settle them: * **Scales must be K-block major.** A block group's sixteen lanes want sixteen consecutive rows of one K block. K-block major coalesces them; the quantiser's own row-major `[M, K/32]` spreads them K/32 bytes apart -- same instruction count, 3.2x the cache accesses, +50-67% on the GEMM. * **Four M tiles pack into one scale dword.** `opsel_b` on the scaled MFMA is an atom-time attribute naming which byte of a 32-bit scale operand to read, so one load serves four tiles and the select is free. -29.1% `SQ_INSTS_VMEM` with VALU flat, worth -5 to -7.5%. A quantiser can emit the layout directly, at no extra traffic. `_PinnedLaunch` resolves each compiled kernel's dispatch once. `JitFunction` re-derived its cache key per call -- an inspect bind, a 35-global snapshot, a drift check -- which measured 181.6us of Python per fused call against 5.9us of actual `hipModuleLaunchKernel`, enough to make the layer host-bound. It verifies the argument tuple element by element before pinning and falls back to the normal path on any surprise, so a FlyDSL change costs performance, not correctness. The fp8 all-gather leg gets its two knobs settled by measurement: pull it over LSA rather than pushing over SDMA (2.0-4.0%), and do not fuse the quantise into the reduce (1.2-1.9%). Both hold at every M and in all three modes. `_gemm_a8w8_8wave.py` is vendored from aiter and now diverges in two places rather than one: `Mfma16x16x128` grew the scaled form of the atom. Both new arguments default to off and the unscaled path is byte-for-byte what it was, so a re-vendor stays a merge. The file's header says so. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * gemm_ar: the GEMM without a collective, for prefill M and for decode's Most of what this integration is worth turns out to be the multiply, not the overlap. At V4.1-Flash's shapes the GEMM is ~180us against 1130-1410us of communication, so hiding it entirely would be worth 11-18% and fusing collects about half of that -- while the GEMM alone is worth ~30% against the route SGLang would otherwise take. And a `ColumnParallelLinear` has no all-reduce to fuse with at all, so `wq_b` was unreachable. `Mxfp8GemmOp` (`gemm.py`) is the same kernel with the epilogue tail compiled out: no communicator, no window, an ordinary tensor out. Two things about it are not obvious. M only has to be a multiple of 64, not of BLOCK_M, because the grid is ceildiv and the tail block masks -- 64 is the packed scale's group. And there is exactly one compile per (N, K) per tile width, because `c_m` is a runtime argument; a per-M cache looks harmless until a server drives it, where every prefill batch pays a fresh multi-second FlyDSL compile. Which N tile to use is decided by the resulting **grid**, not by M. The tile exists to double the grid when the wide one is launch-starved, and the grid is `ceildiv(M,256) * (N/256)`, so a threshold on M alone is only ever right for the N it was fitted to. Across 110 measured points the separation is total: below 128 workgroups the 128-wide tile wins every single time by 9-17%, from 129 the 256-wide one wins every single time by 26-29%. `_WIDE_TILE_MIN_GRID = 140` sits in the gap at zero regret. `Mxfp8GemvOp` (`gemv.py`, `kernels_gemv.py`) is a different kernel for decode's token counts, where a 256-row tile has nothing to fill it. It inverts the same scaled MFMA -- the weight is the A operand and the tokens are B, because the weight is what there is a lot of -- streams the weight from global with no LDS staging, and splits K across a workgroup's waves, reducing through LDS in a fixed order so a row stays batch-invariant. Against SGLang's `mxfp8_gemv` on the same fp8 bytes it is 3-8% faster on the two tuned shapes and **bit-identical** at every M, which is the check that matters: same instruction, same operands, so a layout disagreement shows up as a mismatch rather than a tolerance. Its margin is tuning, and that is recorded next to the table: shapes falling back to the heuristic land inside +-2%, except `wo_a`, which loses 12-16%. The operands are deliberately not uniform. Both ops take `preshuffle_b`'s weight, so a server shuffles once, but the GEMM wants the A scale through `preshuffle_a_scale` while the GEMV takes both scales exactly as the checkpoint stores them. That is not an inconsistency: the GEMM's sixteen lanes read sixteen rows of one K block and have to coalesce, where the GEMV's are sixteen tokens, M is at most 32, and the whole A scale is under a kilobyte. Tests cover both ops at both real shapes, the ragged and unpadded M contracts, and the GEMV across its configuration space -- each configuration in its own process, because a FlyDSL compile failure takes the interpreter with it. Every masked address in the GEMV is pushed out of its buffer's records or clamped: 0xFF is NaN in both ue8m0 and e4m3, and NaN times a zeroed weight is still NaN. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * benchmark, docs: one timer, three entry points, and the traps that cost most The measurements, and the harness that has to be right for them to mean anything. Most of the work here was finding out that it was not. **Three measurement traps, each of which produced a confident wrong number.** A single-call CUDA-graph replay has a ~13.4us floor on this box, which is most of any small-M measurement -- `timing.py` amortises the capture. A repeated call reads the weight out of the 256MB LLC at 1.7x the bandwidth a forward pass gets, so cold is the number that predicts a server -- it rotates copies past the cache. And getting *that* right took two more corrections: the ring advances at capture time, so a graph of `reps` calls only ever touches `reps` copies; and the ring has to be sized so each tensor clears the cache on its own, not so their sum does, because a caller may hand over several weights of which the kernel reads one. Each is documented where it bit, because the symptom is always a plausible number rather than a failure. The worst of them was in the comparison itself: the baseline's *bf16* weight was never rotated, while `native_route_plan` picks `hipblaslt_bf16` -- which reads exactly that tensor -- on most large-M points. Those rows measured a hot baseline against a cold mori, which ran against mori, not for it. Corrected, SGLang is 3.6-18.6% slower there and mori's margin moves up to 28 points in its favour. The rows whose route never reads that tensor are unchanged within a point, which is how the fix was confirmed to touch only what it should. **Failure has to reach the caller.** A correctness sweep that printed `fail=N` and exited 0 -- verified with a probe rigged to fail all 112 cases -- is now a pytest grid. Every benchmark returns non-zero, `sweep.py` exits with the count of failed points, and "mori declines this shape by construction" is recorded as a result rather than a crash, so ten expected declines cannot bury one real break. **Six entry points, not seventeen.** `bench_gemm.py` covers the two scopes that are genuinely different questions -- `--scope kernel` for pre-quantised operands and `--scope linear` for a whole layer -- where three scripts had encoded the distinction implicitly. Six shell sweeps and a tuning script become `sweep.py` presets; two reporters become `report.py`, which keys on every axis a sweep declares varying and so cannot silently overwrite one configuration with another. The SGLang baseline is optional throughout: mori's own numbers need nothing but mori. One historical experiment is kept and renamed: `fused-blockscale-control` ran V4.1-Flash's shapes under a quantisation the checkpoint does not use, while presenting itself as a V4.1 benchmark. `--mode split-lsa --gather-dtype fp8` is refused where it is asked for. The LSA all-reduce has no fp8 leg, and the combination used to run a bf16 collective and report it under the fp8 label; the validation gate's two-sided band caught it, but only after paying for the run. `docs/` keeps how-to-run and a summary; every full table lives in the operator README, next to the code it is about. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * docs: re-measure the end-to-end numbers on AMD's own V4.1 branch The SGLang-side numbers were taken against a synthetic snapshot of the container image, because the V4.1 code existed on no branch at the time. It does now -- sgl-project/sglang#39857, `kevin-mii:dsv41-amd-main` -- so all three servers were re-run on top of it. The fused wo_b reproduces and is slightly better: fp8 wire +5.6% to +6.0% at bs 4-16 against +5.1% to +5.7% before, bf16 +1.9% to +2.8% against +2.3% to +2.6%. GSM8K 0.917 / 0.912 / 0.918. Added, because it had never been in this section: the standalone GEMM on its own, +2.2% to +2.5% with nothing fused, GSM8K 0.920 / 0.928. And a note that the two results do not add -- `fused_wo_b` runs first and the linear only sees what it declines, so they split `wo_b` rather than stacking on it. Both together is a fourth configuration nobody has measured. **One claim here was wrong and is corrected.** It called bs=1 a built-in control on the grounds that M=4096 is below both floors. The floors are 8192 for bf16 and 2048 for fp8, and the shape log confirms the consequence: the bf16 wire's smallest served M is 7734, the fp8 wire's is 1792. bs=1 is a control for the bf16 column only, and the fp8 column's +3.4% there is a real win that was being dismissed as drift. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * docs: the fourth configuration, and one fault that has not reproduced Both SGLang-side paths enabled together, which this section had been calling unmeasured: +7.4% to +7.7% prefill throughput at bs 4-16, against +6.0% for the fp8 wire alone and +2.3% for the standalone GEMM alone. They stack without adding, and the structure says why -- `fused_wo_b` runs first and the linear only sees what it declines, so at `wo_b` they divide the layer by M, and the ~1.7% on top is the column-parallel layers fusing cannot reach. Also recorded: one server out of nine hit `hipErrorIllegalAddress` under sustained load, and re-running the same variant completed cleanly. Written as what it is -- one sample, cause not established, the combination merely the only configuration that has shown it -- rather than as a known defect or as nothing. The five hypotheses already eliminated are listed so the next person does not re-walk them, and so is the reason `AMD_SERIALIZE_KERNEL=3` is being held back: it is worth spending on a repro that reproduces, and serialising perturbs the timing a race would depend on. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * docs: absolute URLs for the links out of the docs tree, as #688 established Rebasing onto main brought in #688, which converted this file's one link into `python/` to an absolute GitHub URL: Sphinx builds `docs/` alone, so a relative path to a file outside that tree is a broken reference and `-W` fails the build. The rewrite in this branch added three more of them. Same rule applied to all four. Links the other way -- from the operator README into `docs/` -- stay relative, because that file is not in the toctree and GitHub resolves them. Verified with the gate itself rather than by eye: `sphinx -E -n -W --keep-going` builds clean, and `tools/check_docs_links.py` reports 8872 local links across 25 pages with 0 errors. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * gemm_ar: fix the CI break, and satisfy the pre-commit hooks Two unrelated CI failures on this branch. They are in one commit because `black` reflowed lines adjacent to the real fix, so splitting them would mean hunk surgery on the same few lines rather than a cleaner history. **The CCO unit test.** All ten `test_standalone_mxfp8_gemv` cases failed with AttributeError: 'Int32' object has no attribute 'type'. Did you mean: 'dtype'? _arith_ops_gen.py:3163, in ShRUIOp.__init__ while the same tests passed here. The generated MLIR builders infer their result type as `operands[0].type`, so they need an operand that is already an `ArithValue`; `fx.Int32` is a wrapper and carries `dtype`. FlyDSL's operators go through `_make_binop` -> `_extract_arith`, which unwraps first, so `(v >> shift) & 0xFF` is right on any FlyDSL that has the operators at all. Whether the wrapper exposes `.type` differs between the FlyDSL here and the one in CI's container, which is why this was invisible locally. The rule that holds on both: hand wrappers to FlyDSL operators, never to raw builders. `kernels_fused._load_byte` carried the identical expression and the identical latent break; CI never reached it because the GEMV tests failed first. The one `arith.andi` over two predicates became nested `arith.select` -- same result, and `select` is what the rest of mori builds with, so it is the op proven against whatever FlyDSL CI ships. Every `arith.*` call left in this branch is now one `origin/main` also uses. 146 tests pass locally, including the fp32-reference numerics that would catch any change in what the shift and mask compute. **pre-commit.** Formatting from the hooks themselves, plus two they flagged and could not fix: - `needs_sglang` in bench_gemm.py was dead (F841). It computed a condition nothing read; `build_mxfp8(want_sglang=...)` already gates the only thing that needs it, and `--scope linear` imports SGLang's quantiser inside `mori_linear_call`, where a missing build is a reported per-impl failure and a non-zero exit. Deleted rather than wired up: there was no guard to restore. - `l` as a comprehension variable in sweep.py (E741) -> `line`. The license hook appended the full MIT header to timing.py above its SPDX short form, leaving two notices; dropped the short one. The second notice in `_gemm_a8w8_8wave.py` is the deliberate vendored-from-aiter attribution and is left alone. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * gemm_ar: close the four gaps from review on the new public surface All four reproduce as reported; none are in the MXFP8 numerics. **`supports_gemm` said yes to a tile the store cannot write.** The permlane guard was `BLOCK_N >= 256`, but `_PermlaneStoreC._emit` pairs *exactly* two N-tiles and asserts `n_tiles_b == 2`. So `block_n=512` passed every predicate and died several frames deeper -- and because the wide tile is only chosen once the grid is large enough, such an instance serves small batches correctly and fails when one grows. That is the worst shape a failure can take. `GemmAllReduceOp` already rejected it at its constructor; `Mxfp8GemmOp` and `supports_gemm` are new here and did not. Rather than restate the rule a third time, it is now `permlane_tile_constraint()` next to `_tile_constraints` in op.py, used by all three, with the kernel's own assert tightened to an equality as the backstop for callers that compile directly. Review suggested the assert alone would cover the predicate; it does not -- `supports_gemm` has to answer without compiling, which is why the rule has to exist as a predicate too. **The benchmarks reported `validated=true` for runs that validated nothing.** `bench_gemm.py` seeded `ref` from the first implementation's output, so that impl was its own reference: `rel_l2` identically 0 and a pass regardless of what it computed. A single-impl run -- which the docs recommend -- checked nothing, and a bug shared by every impl was invisible. `bench_gemv.py` had the same hole via `--baseline none`. Both now score against `reference_partial()`, the fp32 reference `bench_gemm_ar.py` next door already uses, built from the raw operands. The first impl's relL2 is now 1.66e-03 rather than 0, which is the fp8 floor and matches what review measured independently. Where no independent reference exists -- anything quantising inside the timed region -- `validated` is `None` and prints as `n/a`, never `True`. **`--scope kernel --impl sglang` compared two different inputs.** mori's kernel path builds its own pre-quantised A from seed 99 while SGLang's entry point is a linear and quantises the outer bf16 `x`. The baseline was scored against an unrelated result and read as failing validation -- a false accusation -- and the timings were not comparable either, only one side carrying the quantisation. Now refused at argument parsing, which is what this file's own docstring already said: linear is the scope where that comparison means anything. **`sweep.py fused-fp8` shipped 3 points that cannot run.** The preset listed `split-lsa` while appending `--gather-dtype fp8`, and the bench refuses that pair at entry -- a rejection this branch added and did not propagate here. 15 points, 3 of them unrunnable on a healthy box. Dropped from this preset only; the bf16 preset keeps `split-lsa`, where it is legal. Not addressed: the `wo_a` tuning entries review found worth 13.9-17.5%. That is a new performance claim and wants its own reproduction rather than being taken on the review's numbers inside an already-large PR. 66 tests pass; both kernel-scope paths verified against the new reference. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> --------- Co-authored-by: xiangch <xiangch@amd.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
5 tasks done
…top-k epilogue gfx950 ports of the DeepSeek V4 JIT kernels (index-K/Q fp4 packing, compressor stores, mHC boundary, MXFP8 GEMV, fused K norm-rope) with the native fp8 type aliases, and an in-kernel sort of each row's picks in the AOT top-k transform (sort_output) so the sparse kernels sum in a fixed order. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
…n ROCm The fp4 indexer host (one-launch index-Q pack, sorted paged top-k, candidate blocks), the RoPE fake-quant kernels, the mHC boundary with the Sinkhorn reduce hosted in the next RMSNorm launch, and the decode glue launches (length fold, compression metadata, page table) that replace the torch glue on the launch-bound decode step. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
The ROCm router gate (split-K GEMV up to 64 rows, one-launch gate + sort at M <= 2), the shared-expert add folded into aiter's FlyDSL top-k reduction, the gfx950 native MXFP8 route and fp8-grid producers, and the split-K bf16 x fp8 batched GEMM for the wo_a decode projection. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
… aiter sparse decode The HIP radix attention backend (breakable prefill graphs, DSpark verify metadata, aiter_sparse decode with a pinned split-KV count, decoder SWA bounded replay), the FlyDSL low-ratio indexer, FlashMLA glue, and the aiter fp4 index-K cache layout behind one hook in the CUDA backend and the memory pool. The decode top-k is ordered by position so attention bits do not follow the KV page layout. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
aiter MoE runner with fused sorting and the fused reduce-add scope, the one-launch MoE front, gfx950 fp8 dispatch into fp8_hip, the clamped SiLU, and the fused reduce-add gate that tolerates a fused or absent shared expert. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
…dense route, Engram if _is_hip dispatches from the model into deepseek_common/amd (fused mHC boundary with deferred coefficients, gfx950 dense helpers, index-Q one-launch, fused decode glue, SWA tail slicing), the Engram image select, the new SGLANG_OPT_HIP_* switches, and the ROCm defaults for the aiter batched GEMM and the FlashMLA backend. The Engram hash kernel runs on HIP too, so upstream's in-kernel decode commit (sgl-project#39138) serves both platforms and the history is written once per step. deepseek_v4_low_ratio_sources and _every_row_routed follow the ForwardBatch rename num_token_non_padded_cpu -> global_num_token_non_padded_cpu (the old name is no longer a field), which the HIP low-ratio path reads under the breakable prefill graph. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
With --enable-dp-attention every DP rank must replay the same captured decode graph: the in-graph collectives pair the ranks' registered buffers by capture order. The variant was chosen from the rank's own lengths, so an idle rank replayed a different graph from the decoding rank and the busy rank read a never-written peer buffer (garbage from position 1 on MI350X DP2). Carry the group max seq len through the scheduler's existing per-step DP sync and select from it. Verified on HIP; Hopper captures no variants. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
…d FMoE rows aiter 4ad99832 with two cherry-picks and a patch that drains the LDS-DMA group before the stage-1 K-step barrier (a8w4 stage 1 was not bitwise repeatable without it), the DeepSeek-V4.1 EP4 a8w4 FMoE rows merged through aiter's model_configs glob, the documented SGLANG_OPT_HIP_* switches, and a pre-commit exclude so the whitespace fixers leave patch files intact. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
Bitwise and tolerance guards for every HIP kernel and path the MI350X recipes take, plus regression tests for the bugs fixed here (fold-cache invalidation, decode selection order, sorted top-k padding, -1 keys, saturation and tie rounding, the fused reduce gate, split-K row count, breakable-graph refresh contract). Registered for the MI35x suite with measured est_time values. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_019tZvxNDYPWtWc7ZtfXuybA
Keep the AITER synchronization fix, but leave global ROCm defaults and model-specific FMoE tuning to a separate deployment change.
Use one context line per hunk so the patch needs no whitespace-hook exclusions. Preserve both synchronization fixes and shorten their comments.
Keep the existing MI35x CI suites while conforming to the target branch registered-test layout. Test contents and runner selection are unchanged.
kevin-mii
requested review from
Alisehen,
DarkSharpness,
HydraQYH,
JustinTong0323,
Kangyan-Zhou,
OrangeRedeng,
b8zhong,
bingxche,
celve,
hebiao064,
ishandhanani,
mmangkad,
sogalin,
wisclmy0611,
yctseng0211,
yuan-luo and
zijiexia
as code owners
September 23, 2026 23:38
…ethod The encoder-replay floor stays computed inline as main does: window_layout takes no replay flag, so the call raised TypeError whenever a request window was set. The HIP prelaunch of the compressor KV score called a method no compressor defines. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
dp_max_seq_len (scheduler DP sync, ScheduleBatch, ForwardBatch, the graph variant policies) and the single-rank DP check for decoder-tail replay change CUDA behavior too; they now live on dp-graph-variant-sync off main. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Restores main's wording, backticks and :func: references where only the text changed, main's q_for_wqb name and constant order, and the cookbook. Two comments change minimally because main's text would now be false: the aiter batched GEMM default and the V4.1 layout support list. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The branch wrapped aiter.fused_moe.moe_sorting (a one-launch small-M sort, also folded into the ROCm router gate) and FlyDSL's _run_moe_reduction (adding the shared expert in the reduction). Without them aiter's kernels run and select_experts masks the padded rows as on main; serving is unchanged within noise (real b1 452 vs 457, b8 1591 vs 1582 tok/s). Removes the three kernels, the router's pending-sort handoff, the deferred-pad top-k output and the shared-add hooks in deepseek_v2 and the MXFP4 helper; aiter.py and its runner test are main's again. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
1 task done
Measured on MI355X (DeepSeek-V4.1-Flash TP4, DSpark EP4) with AITER_BF16_FP8_MOE_BOUND=0, as the AMD DeepSeek-V4 tests and cookbook set it: - the stage-1 LDS-DMA drain: GSM8K 0.885 unpatched vs 0.905 patched (200 questions), throughput within noise, and its race test passes unpatched; it stays upstream as ROCm/aiter#5561; - the bf16 SiLU route is never taken at bound 0; the fix is ROCm/aiter#5802; - the tuned FMoE CSV shows no end-to-end gain; it stays upstream as ROCm/aiter#5562. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
This was referenced Sep 24, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Run DeepSeek-V4.1-Flash on AMD gfx950 with HIP attention, quantization, MoE and DSpark paths. This continues the AMD work from #39186 on a branch targeting
main.Dependency stack: main now includes the compression/KV metadata extraction (#39652). This branch still carries the remaining V4.1 prerequisites from
dsv4.1, including #38798 and #39666, which are open. Those shared model, vision/chat, CUDA and Engram changes explain the larger main-relative diff. Keep this PR draft until those prerequisites land, then rebase to expose the AMD-only diff. They should not be reviewed as newly authored AMD ports.Modifications
Accuracy Tests
On MI350X, ROCm 7.2.4, with isolated patched AITER:
f5cd9c791c: AMD, main prefill-graph, DP scheduling and model-guard suite 203 passed, 40 skipped, 302 subtests. All-files pre-commit passed.0171bde4f2: 203 tests and 302 subtests passed, 40 skipped. Focused quantization/KV/C2/WO-A tests passed on all eight GPUs. Eight-rank entrypoints passed: DSpark 36 tests/36 NVLink-only skips, Engram 5 tests, all-reduce/mHC 35 tests (rank-zero report counts, all ranks exited successfully).Serving used
deepseek-ai/DeepSeek-V4.1-Flashrevisiondba1be0a40aa45a94ad051997016db3960a90277, FP8 KV, page size 256, radix cache disabled, shared-expert fusion disabled, DSpark block size 5. Environment:TRITON_HIP_USE_ASYNC_COPY=0,SGLANG_USE_AITER=1,SGLANG_USE_ROCM700A=0,ROCM_QUICK_REDUCE_QUANTIZATION=NONE,AITER_BF16_FP8_MOE_BOUND=0.DSpark-off serving was measured at
629d905531; DSpark-on at0171bde4f2after a raw-verification metadata argument fix. Model/kernel executable ASTs/tokens remain equivalent on the publication branch; the fresh regression results above supplement those serving runs.Speed Tests and Profiling
Four MI350X GPUs, TP4/EP4, 4,096 input / 1,024 output tokens,
ignore_eos, six timed runs after a discarded warm-up. Median output tokens/s from first to last streamed event, cache flushed per run. Real DSpark acceptance; no acceptance simulation.These are overall DSpark benefits, not incremental gains attributable to this PR. Random-input DSpark throughput was unchanged across the rebase: 641.89 → 641.88 tokens/s. Real-text changes versus the prior revision were +2.64%, +2.65%, and −1.15%, with acceptance lengths also changing. One launch per mode; this is not an interleaved noise-controlled comparison.
The random prompt is reconstructed at seed 42; the original blog
prompt.jsonwas unavailable. The original PR's 548–556 tokens/s used simulated acceptance fixed at 5.5 and is not directly comparable to these real-acceptance measurements.Remaining validation and follow-ups
Not executed: CUDA or ROCm 10 validation, a fresh Docker build, and full-model CP/EAGLE configurations blocked by existing guards. CUDA-only small-metadata reference tests require
sgl_kernel.flashmla_ops, unavailable in the ROCm environment. Successful collective tests do not establish that the previously intermittent RCCL graph-capture abort is fixed.The independent AITER tuning comparisons were operator benchmarks; their numerical tolerances still need independent validation. The synchronization regression exercises the path but does not reproduce the historical intermittent race. No new end-to-end claim is inferred from those operator results.
Checklist
CI States
Latest PR Test (Base): ❌ Run #35944357239
Latest PR Test (Extra): ❌ Run #35944356709
Latest PR Test (AMD ROCm 10): ❌ Run #35944357043