Skip to content

[RDNA2] Port rdna_extras onto vLLM v0.31.0 (rdna_extra/v0.31.0) - #47

Merged
BlivionIaG merged 170 commits into
rdna_extrasfrom
rdna_extra/v0.31.0
Oct 11, 2026
Merged

[RDNA2] Port rdna_extras onto vLLM v0.31.0 (rdna_extra/v0.31.0)#47
BlivionIaG merged 170 commits into
rdna_extrasfrom
rdna_extra/v0.31.0

Conversation

@BlivionIaG

@BlivionIaG BlivionIaG commented Oct 7, 2026 •
edited
Loading

Copy link
Copy Markdown

Ports the RDNA fork (c6d99ca305, based on v0.28.0) onto upstream vLLM v0.31.0. The branch is rdna_extra/v0.31.0 = v0.31.0 + 170 topic commits.

Branch layout (2026-10-11): the old 0.28 line is now rdna_extra/v0.28.0 (renamed from rdna_extras; the open 0.28-era PRs moved with it), and rdna_extras is upstream v0.31.0 (db9527a468), the repo's default branch and this PR's base. Merging this PR makes rdna_extras the v0.31 RDNA line.

✅ Validated on 20edef5321 (4× / 2× Radeon PRO V620, par1-cs25): every serving config passes all gates (greedy probe before/after load, PPL, prefix cache, concurrency, 0 GPU faults): Flash-Next MTP0 / MTP2 / agentic, 27B AWQ (+W4A8, +TP=2 full.env), Qwen3.6-35B-A3B MTP0 / MTP2, DSV4-Flash. Unit suite: only known pre-existing failures. Final numbers in Validation status → Final pass.

The PR diff now shows only the port (base rdna_extras = v0.31.0): git log v0.31.0..rdna_extra/v0.31.0 (170 commits).

⚡ Opt-in W4A8 prefill (W4A8=1 / VLLM_RDNA2_W4A8_SDOT4=1): on Qwen3.8-27B AWQ it gives +24-34 % prefill (16k TTFT 18.5 → 14.9 s) with byte-identical greedy output and decode at parity. Off by default; recommended for prefill-heavy serving. Details: port notes.

How it was built

  • I merged the fork's net delta onto v0.31.0 in one pass, using the 3-way base e25c586b90, the last upstream main commit the fork had merged. That gave 102 conflicts, all resolved by hand. I then split the result into commits by topic and by each file's dominant original author. Intermediate commits may not build; the tip does.
  • Authors and co-authors come from the original commits. Each commit lists the original SHAs it carries. Agent identities map to BlivionIaG.
  • qwen4_exp: upstream merged its own Qwen3.8 Flash-Next ([Model] Support Qwen3.8-Flash-Next vllm-project/vllm#53896). The fork's copy came from an earlier revision of that PR, so I used upstream's e126687a9a as the merge base and replayed only the real RDNA deltas.

Decisions and open questions (full list in docs/rdna2/v0.31.0-port-notes.md)

  1. PLE offload: upstream's by default. EngramConfig.cpu_offload (pinned-host UVA, [Qwen4Exp][ROCm] PLE n-gram table CPU offload vllm-project/vllm#57497) is used, with our fixes to add on top after testing.
    • The Flash-Next recipes no longer use the CPU sidecar.
    • The sidecar code stays in the tree but is not wired in; VLLM_PLE_CPU_OFFLOAD=1 fails with an explicit error.
  2. Qwen4Exp KV grouping: nothing to carry over. The fork's CSA+linear grouping came from a pre-merge revision of upstream's own [Model] Support Qwen3.8-Flash-Next vllm-project/vllm#53896, and upstream's packed KV groups replace it. Only KV memory and prefix caching need re-checking.
  3. QSA prefix-cache ring fix (e45dd5cb2) is replaced by upstream's scratch-group skipping. Prefix hits must be re-checked.
  4. QSA prefill bound (fork PR perf(rocm): restore V620 startup and bound QSA prefill work #15, +18 %) is replaced by upstream's fields and workspace changes. Needs a re-measure.
  5. Launchers moved to tools/rdna/ (upstream already owns tools/recipes/). Run them with bash tools/rdna/serve_rdna.sh RECIPE=….
  6. Upstream behaviour off RDNA; fork paths gated to ROCm RDNA. The gate is vllm/platforms/rdna.py::on_rdna_family() (gfx10x or RDNA3/4).
    • Affected areas: TP collectives, custom all-reduce, compile backend, breakable CG, GDN layer and metadata, norms and RoPE, V1/V2 runner FULL handling, scheduler.
    • Where upstream already tunes RDNA3/4, those GPUs keep upstream and the fork change applies to gfx10x only.
    • Fixed: the EXL3 fp16 cast never fired on gfx1030, because upstream on_rdna() excludes gfx10x.

Superseded by upstream (fork code dropped)

  • FULL_AND_PIECEWISE → FULL_DECODE_ONLY downgrade
  • MTP under PP (broadcast_drafts)
  • kv-offload scratch groups and MTP fixes
  • replay boundaries
  • slot_mapping_enabled
  • get_mamba_state_copy_funcs
  • DeepSeek MTP SupportsPP

Pre-existing bugs fixed in passing

  • RDNA2 W8A16 MoE method: missing set_weight_attrs import, which raised NameError at weight creation.
  • test_exl3_rdna2.py: missing numpy import.
  • ple_offload/worker.py: ported off initialize_dummy_weights, which upstream removed.

What was verified (no GPU)

At port time, before any GPU run:

  • No conflict markers left. All 225 changed Python files parse.
  • ruff F821/F811/F601/F7/F63 checks are clean.
  • A custom AST pass looked for duplicate class fields, methods and kwargs. It found four silent duplicates from the auto-merge, two of them syntax errors; all are fixed.
  • Every vllm.* import in the changed files resolves against the v0.31 tree.
  • No duplicate op names in csrc/rocm/torch_bindings.cpp.
  • Build, unit tests, pre-commit and mypy came later on the box (see the test plan).

Validation status (updated 2026-10-10)

Drivers: tools/rdna/port_v031/{build,unit,serve_validate}.sh. Raw results are in ~/w4a8_runs/port-v031/ on the box; the full log is in docs/rdna2/v0.31.0-port-notes.md → Validation log.

Headline metrics (TP=4, 4× V620, F&P cudagraphs)

History only. Current numbers: Final pass on 3b43339ab0 → 20edef5321 below.

1k/512 and 16k/1k at c=1 and c=8. Prefill = mean input / mean TTFT per request; decode = 1000 / mean TPOT per request. Flash-Next and 27B AWQ (v0.31) rows are the re-validation with the rdna_ar fix (00fcc97ff5, 2026-10-08 21:34); all passed probe, prefix and after-load probe, 0 faults. EXL3 rows predate the fix.

Model Cell Prefill tok/s Decode tok/s/req Agg out tok/s TTFT s ITL med ms
Flash-Next W4A16 MTP0 (v0.31) 1024/512 c=1 1526 41.1 39.1 0.67 24.3
Flash-Next W4A16 MTP0 (v0.31) 1024/512 c=8 494 24.0 175.3 2.07 37.9
Flash-Next W4A16 MTP0 (v0.31) 16384/1024 c=1 1972 40.6 30.6 8.31 24.6
Flash-Next W4A16 MTP0 (v0.31) 16384/1024 c=8 917 12.5 81.6 17.88 38.3
Flash-Next W4A16 MTP2 (v0.31) 1024/512 c=1 1963 70.4 65.8 0.52 31.4
Flash-Next W4A16 MTP2 (v0.31) 1024/512 c=8 898 25.7 188.1 1.14 69.0
Flash-Next W4A16 MTP2 (v0.31) 16384/1024 c=1 1942 72.1 45.3 8.44 31.8
Flash-Next W4A16 MTP2 (v0.31) 16384/1024 c=8 1146 13.6 88.3 14.30 69.7
Qwen3.8-27B AWQ (v0.31) 1024/512 c=1 1070 39.9 37.2 0.96 25.1
Qwen3.8-27B AWQ (v0.31) 1024/512 c=8 274 21.6 149.3 3.73 39.0
Qwen3.8-27B AWQ (v0.31) 16384/1024 c=1 884 35.5 21.6 18.54 28.2
Qwen3.8-27B AWQ (v0.31) 16384/1024 c=8 476 6.8 44.0 34.46 55.6
Qwen3.8-27B EXL3 MTP0 (v0.31, 36c4164b22) 1024/512 c=1 1221 20.4 19.8 0.84 48.9
Qwen3.8-27B EXL3 MTP0 (v0.31, 36c4164b22) 1024/512 c=8 275 11.9 87.6 3.73 78.1
Qwen3.8-27B EXL3 MTP0 (v0.31, 36c4164b22) 16384/1024 c=1 1019 19.3 14.8 16.08 51.9
Qwen3.8-27B EXL3 MTP0 (v0.31, 36c4164b22) 16384/1024 c=8 473 5.5 36.9 34.66 92.8
Qwen3.8-27B EXL3 MTP2 greedy (v0.31, 36c4164b22) 1024/512 c=1 1111 25.9 24.8 0.92 113.8
Qwen3.8-27B EXL3 MTP2 greedy (v0.31, 36c4164b22) 1024/512 c=8 323 6.0 45.1 3.17 465.0
Qwen3.8-27B EXL3 MTP2 greedy (v0.31, 36c4164b22) 16384/1024 c=1 931 23.9 17.0 17.60 124.2
Qwen3.8-27B EXL3 MTP2 greedy (v0.31, 36c4164b22) 16384/1024 c=8 470 4.3 29.5 34.84 497.1
Qwen3.6-35B-A3B GPTQ MTP0 (v0.31, eec301a25a) 1024/512 c=1 5193 77.6 75.5 0.2 12.9
Qwen3.6-35B-A3B GPTQ MTP0 (v0.31, eec301a25a) 1024/512 c=8 1354 51.9 386.1 0.76 18.1
Qwen3.6-35B-A3B GPTQ MTP0 (v0.31, eec301a25a) 16384/1024 c=1 4153 68.1 54.0 3.95 14.7
Qwen3.6-35B-A3B GPTQ MTP0 (v0.31, eec301a25a) 16384/1024 c=8 1747 21.0 140.8 9.38 26.5
Qwen3.6-35B-A3B GPTQ MTP2 greedy (v0.31, eec301a25a) 1024/512 c=1 5131 160.6 151.4 0.2 16.6
Qwen3.6-35B-A3B GPTQ MTP2 greedy (v0.31, eec301a25a) 1024/512 c=8 968 56.0 391.3 1.06 43.9
Qwen3.6-35B-A3B GPTQ MTP2 greedy (v0.31, eec301a25a) 16384/1024 c=1 3829 135.3 86.5 4.28 20.8
Qwen3.6-35B-A3B GPTQ MTP2 greedy (v0.31, eec301a25a) 16384/1024 c=8 1933 22.6 150.3 8.48 60.5
DeepSeek-V4-Flash K128 fp16 (v0.31, 092d7ea70c, SEQS=8) 1024/512 c=1 1082 38.8 36.2 0.95 25.8
DeepSeek-V4-Flash K128 fp16 (v0.31, 092d7ea70c, SEQS=8) 1024/512 c=8 286 22.0 152.5 3.58 40.1
DeepSeek-V4-Flash K128 fp16 (v0.31, 092d7ea70c, SEQS=8) 16384/1024 c=1 1092 37.3 24.1 15.0 26.8
DeepSeek-V4-Flash K128 fp16 (v0.31, 092d7ea70c, SEQS=8) 16384/1024 c=8 564 8.6 55.0 29.08 43.8
Qwen3.8-27B AWQ W4A8 on (v0.31, b03bf9d210) 1024/512 c=1 1417 43.3 40.9 0.72 23.1
Qwen3.8-27B AWQ W4A8 on (v0.31, b03bf9d210) 1024/512 c=8 335 23.6 165.8 3.05 37.3
Qwen3.8-27B AWQ W4A8 on (v0.31, b03bf9d210) 16384/1024 c=1 1102 38.2 24.6 14.86 26.2
Qwen3.8-27B AWQ W4A8 on (v0.31, b03bf9d210) 16384/1024 c=8 496 7.4 47.5 33.03 54.0
Qwen3.8-27B AWQ W4A8 off, same run (v0.31, b03bf9d210) 1024/512 c=1 1113 42.1 39.2 0.92 23.7
Qwen3.8-27B AWQ W4A8 off, same run (v0.31, b03bf9d210) 1024/512 c=8 250 22.1 150.4 4.1 38.6
Qwen3.8-27B AWQ W4A8 off, same run (v0.31, b03bf9d210) 16384/1024 c=1 885 37.3 22.3 18.51 26.8
Qwen3.8-27B AWQ W4A8 off, same run (v0.31, b03bf9d210) 16384/1024 c=8 400 6.4 40.5 40.95 55.0
Qwen3.8-27B EXL3 MTP2 greedy (v0.31, f419d8d6c0) 1024/512 c=1 1211 80.1 70.9 0.85 36.6
Qwen3.8-27B EXL3 MTP2 greedy (v0.31, f419d8d6c0) 1024/512 c=8 420 29.6 200.1 2.44 70.1
Qwen3.8-27B EXL3 MTP2 greedy (v0.31, f419d8d6c0) 16384/1024 c=1 949 64.3 30.9 17.3 46.1
Qwen3.8-27B EXL3 MTP2 greedy (v0.31, f419d8d6c0) 16384/1024 c=8 486 8.5 53.0 33.7 101.9
Qwen3.8-27B EXL3 MTP0 (v0.31, f419d8d6c0) 1024/512 c=1 1215 32.9 31.3 0.84 30.4
Qwen3.8-27B EXL3 MTP0 (v0.31, f419d8d6c0) 1024/512 c=8 288 19.8 139.6 3.55 44.4
Qwen3.8-27B EXL3 MTP0 (v0.31, f419d8d6c0) 16384/1024 c=1 1014 30.0 20.4 16.2 33.3
Qwen3.8-27B EXL3 MTP0 (v0.31, f419d8d6c0) 16384/1024 c=8 452 6.7 43.1 36.3 59.4
Qwen3.8-27B AWQ (0.28) 1024/512 c=1 1117 42.2 39.3 0.92 23.7
Qwen3.8-27B AWQ (0.28) 1024/512 c=8 251 22.8 154.5 4.08 37.2
Qwen3.8-27B AWQ (0.28) 16384/1024 c=8 401 6.4 40.9 40.82 53.8
Qwen3.8-27B EXL3 MTP0 (v0.31) 1024/512 c=1 1203 19.1 18.5 0.85 52.4
Qwen3.8-27B EXL3 MTP0 (v0.31) 1024/512 c=8 272 11.4 84.1 3.76 81.7
Qwen3.8-27B EXL3 MTP0 (v0.31) 16384/1024 c=8 465 5.3 36.0 35.20 97.0
Qwen3.8-27B EXL3 MTP2 (v0.31) 1024/512 c=1 1151 22.0 21.2 0.89 119.9
Qwen3.8-27B EXL3 MTP2 (v0.31) 1024/512 c=8 374 4.9 35.8 2.74 473.8
Qwen3.8-27B EXL3 MTP2 (v0.31) 16384/1024 c=8 467 3.5 23.5 35.09 506.2

Final pass on 3b43339ab0 → 20edef5321 (2026-10-10 17:06 → 2026-10-11 03:58)

Same gates as below, fresh cache per run, TP=4 F&P on HIP 6-9, greedy benches. Covers what landed since f419d8d6c0: fused hyper-connection prefill (default on), stall-cap fit, replay-copy wrapper removed. Cells: prefill tok/s/req · decode tok/s/req · TTFT s · ITL median ms · aggregate output tok/s; (previous plan in brackets).

Run Gates 1k/512 c=1 1k/512 c=8 16k/1k c=1 16k/1k c=8
Flash-Next MTP0 all PASS, PPL 4.80, prefix 22.8×, concurrency 0 fail 2006 · 45.0 · 0.51 · 22.2 · 43.1 (1511 · 43.4 · 0.68 · 23.0 · 41.1) 524 · 25.8 · 1.95 · 35.2 · 188.0 (497 · 25.1 · 2.06 · 36.1 · 182.3) 2117 · 44.5 · 7.74 · 22.5 · 33.3 (1979 · 42.9 · 8.28 · 23.3 · 31.9) 1078 · 13.5 · 15.19 · 35.3 · 89.4 (1033 · 12.9 · 15.9 · 36.9 · 85.7)
Flash-Next MTP2 all PASS, PPL 4.80, prefix 13.1×, concurrency 0 fail 2116 · 69.2 · 0.48 · 28.9 · 65.1 (1957 · 68.2 · 0.52 · 29.8 · 63.9) 936 · 29.1 · 1.09 · 60.7 · 204.5 (908 · 24.8 · 1.13 · 72.4 · 172.7) 2052 · 69.0 · 7.99 · 29.1 · 44.9 (1948 · 79.3 · 8.41 · 30.2 · 48.0) 1303 · 14.0 · 12.57 · 61.7 · 92.6 (1114 · 10.6 · 14.7 · 73.8 · 72.1)
  • Flash-Next MTP0: every cell faster than the previous plan (prefill +4-33 %, decode +3-4 %, lower TTFT and ITL). Stall probe at defaults (cap off on Flash-Next, as intended): 16k injection max gap 504 ms, 114 decoder tokens during the prefill; 4k 494 ms; reverse 6.8 s, unchanged from the baseline.
  • Flash-Next MTP2: c=8 aggregate output +18 % (1k) and +28 % (16k), back above the pre-gemv-regression bests (201.0 / 92.8). 16k c=1 decode 69.0 vs 79.3 tok/s/req is acceptance variance (ITL median 29.1 vs 30.2 ms; this cell ranged 72-83 across earlier runs).
    | 27B AWQ (cap on) | all PASS, PPL 4.94, prefix 25.3×, concurrency 0 fail | 1102 · 44.7 · 0.93 · 22.4 · 41.4 (1099 · 44.0 · 0.93 · 22.7 · 40.8) | 374 · 22.0 · 2.74 · 36.3 · 157.5 (256 · 23.1 · 4.01 · 36.6 · 156.9) | 897 · 39.5 · 18.27 · 25.3 · 23.2 (894 · 38.9 · 18.3 · 25.7 · 23.0) | 406 · 6.4 · 40.40 · 52.4 · 40.8 (404 · 6.5 · 40.5 · 52.8 · 41.2) |
    | Flash-Next, capture sizes ≤512, concurrency ×2 | PASS / PASS, 0 failing (the earlier one-off FAIL did not reproduce in 3 repeats: wrapper removal closed) | — | — | — | — |

| Qwen3.6-35B-A3B MTP0 (cap on, MoE fit) | all PASS, PPL 5.49, prefix 13.3×, concurrency 0 fail | 5698 · 77.6 · 0.18 · 12.9 · 75.7 (5041 · 75.2 · 0.20 · 13.3 · 73.2) | 1494 · 52.0 · 0.69 · 18.1 · 389.5 (1433 · 51.1 · 0.71 · 18.4 · 381.8) | 4143 · 68.5 · 3.95 · 14.6 · 54.2 (4002 · 66.4 · 4.09 · 15.1 · 52.5) | 1745 · 20.2 · 9.39 · 26.4 · 135.8 (1670 · 20.6 · 9.81 · 27.0 · 137.2) |
| Qwen3.6-35B-A3B MTP2 (greedy, cap on, MoE fit) | all PASS, PPL 5.46, prefix 6.5×, concurrency 0 fail | 5628 · 163.1 · 0.18 · 16.5 · 154.4 (4995 · 152.0 · 0.21 · 17.1 · 143.5) | 1644 · 69.5 · 0.62 · 32.9 · 500.1 (2113 · 59.1 · 0.48 · 40.9 · 439.9) | 3908 · 135.8 · 4.19 · 20.7 · 87.3 (3750 · 126.1 · 4.37 · 21.5 · 82.0) | 1366 · 23.8 · 11.99 · 50.8 · 147.0 (1686 · 23.4 · 9.72 · 58.0 · 151.3) |
| Flash-Next agentic (flashnext.env, SEQS=6) | all PASS, PPL 4.84, prefix 20.4×, concurrency 0 fail | 1601 · 43.1 · 0.64 · 23.2 · 40.9 (1512 · 41.7 · 0.68 · 24.0 · 39.6) | 137 · 28.9 · 7.47 · 32.1 · 148.1 (135 · 28.4 · 7.60 · 32.5 · 145.5) | 2026 · 42.6 · 8.09 · 23.5 · 31.9 (1971 · 41.2 · 8.31 · 24.3 · 30.9) | 459 · 16.0 · 35.73 · 32.4 · 77.9 (451 · 15.7 · 36.3 · 33.2 · 76.4) |
| 27B AWQ W4A8=1 (cap on) | all PASS, PPL 4.89, prefix 22.7×, concurrency 0 fail | 1441 · 44.9 · 0.71 · 22.3 · 42.4 (1436 · 44.1 · 0.71 · 22.7 · 41.6) | 547 · 22.5 · 1.87 · 36.1 · 165.9 (337 · 24.1 · 3.04 · 36.5 · 168.7) | 1124 · 39.7 · 14.58 · 25.2 · 25.4 (1112 · 39.0 · 14.7 · 25.7 · 25.0) | 458 · 7.2 · 35.81 · 52.1 · 45.7 (500 · 7.5 · 32.7 · 52.8 · 48.1) |
| ❌ full.env 27B AWQ TP=2 | FAIL: greedy probe 0/3 (" is is is…", "====="), concurrency 16 failing, after-load FAIL; PPL 4.94 PASS (prefill fine, decode broken) | 572 · 30.1 · 1.79 · 33.2 · 27.3 | 77 · 18.0 · 13.25 · 44.3 · 90.4 | 505 · 25.8 · 32.47 · 38.8 · 14.2 | 120 · 5.0 · 136.6 · 78.0 · 23.1 |
| DSV4-Flash K128 | all PASS, PPL 8.40, prefix 47.7×, concurrency 0 fail (gates only) | — | — | — | — |
| ✅ full.env 27B AWQ TP=2 on dc3065852b (fix) | all PASS, PPL 4.94, prefix 64.2×, concurrency 48/48 identical | 572 · 30.1 · 1.79 · 33.2 · 27.3 (559 · 29.9 · 1.83 · 33.5 · 27.0) | 74 · 17.6 · 13.89 · 44.3 · 87.9 (61 · 20.2 · 16.8 · 45.0 · 89.2) | 504 · 25.8 · 32.53 · 38.8 · 14.2 (491 · 25.5 · 33.4 · 39.3 · 13.9) | 82 · 5.1 · 200.3 · 211.1 · 19.6 (122 · 4.8 · 134.6 · 79.8 · 22.7) |
| DSV4-Flash K128 on dc3065852b (breakable path) | all PASS, PPL 8.19, prefix 47.5×, concurrency 0 fail | — | — | — | — |
| ✅ Flash-Next agentic on the mtp0 performance block (20edef5321: fused hyper-connection on) | all PASS, PPL 4.81, prefix 23.9×, concurrency 0 fail | 2049 · 44.6 · 0.50 · 22.4 · 42.8 (1601 · 43.1 · 0.64 · 23.2 · 40.9) | 140 · 29.1 · 7.33 · 31.9 · 150.1 (137 · 28.9 · 7.47 · 32.1 · 148.1) | 2126 · 44.1 · 7.70 · 22.7 · 33.1 (2026 · 42.6 · 8.09 · 23.5 · 31.9) | 472 · 16.3 · 34.69 · 32.3 · 79.6 (459 · 16.0 · 35.73 · 32.4 · 77.9) |
| ✅ full.env 27B AWQ TP=2, stall cap off (2d6e1cd184, recipe default now) | all PASS, PPL 4.94, prefix 64.9×, concurrency 0 fail | 574 · 30.2 · 1.78 · 33.2 · 27.3 | 62 · 20.5 · 16.45 · 44.2 · 90.8 | 505 · 25.8 · 32.43 · 38.8 · 14.2 | 127 · 5.0 · 129.1 · 77.8 · 23.6 (cap on: 19.6; previous plan: 22.7) |
| 27B AWQ re-measure on e09a8c9509 (dense → all-steps fit) | all PASS, PPL 4.94, prefix 25.3×, concurrency 0 fail | — | 370 · 22.0 · 2.76 · 36.2 · 157.6 | — | 335 · 6.2 · 48.98 · 52.5 · 38.0 |

  • Gates: 10 of 11 runs PASS, 0 GPU faults. The replay-copy wrapper removal is confirmed on Flash-Next MTP0/MTP2 (incl. the 512-size ladder ×2), 27B AWQ (+W4A8), Qwen3.6 MTP0/MTP2, Flash-Next agentic and DSV4.
  • ❌ full.env (27B AWQ TP=2) decodes garbage; it passed on f419d8d6c0. It is the only recipe that forces breakable cudagraphs on (VLLM_USE_BREAKABLE_CUDAGRAPH=1, TP=2, 16-token blocks, ATTN=none), and the wrapper removal changed the breakable-graph code (weak refs in eager breaks). Cause found: A/B on the same recipe, breakable graphs off → PASS 3/3, concurrency 0 fail; pre-wrapper-removal tree → PASS 3/3, 48/48 identical. The removal switched RDNA eager breaks to upstream weak refs; on RDNA the inductor temps behind GDN/FA inputs are not always graph-pool owned, so the weak refs dangle into capture-time dummy activations. Fixed in dc3065852b (RDNA strong-ref branch restored) and re-validated: full.env all gates PASS (concurrency 48/48 identical), DSV4 gates PASS. Graph memory with RDNA strong refs 0.76 GiB vs 0.73 with weak refs and 1.08 with the old wrapper. The stall cap cost full.env (TP=2) 14 % at 16k c=8 (19.6 vs 22.7 tok/s, TTFT 200 vs 135 s), so it is off in that recipe (2d6e1cd184): re-run 23.6 tok/s, TTFT 129 s, all gates PASS. The cap stays on in the TP=4 dense recipes.
  • Stall cap on 27B, fixed: with the dense all-steps fit (e09a8c9509, log: "fit on all prefill steps") the 16k injection max decode gap is 320 ms (all 3 repeats; 4k: 244 ms), 526 decoder tokens during the 16k prefill, vs 776 ms with the mixed-steps-only fit and 2672 ms without the cap. Cost: 16k c=8 aggregate output 38.0 vs 41.2 tok/s without the cap (−8 %), TTFT 49 s; 1k c=8 unchanged.
  • Cap trade-off also visible on W4A8 (16k c=8 45.7 vs 48.1 tok/s, 1k c=8 TTFT 3.04 → 1.87 s; this run used the old mixed-only fit) and on Qwen3.6 MTP2 16k c=8 (147.0 vs 151.3; MoE fit, unchanged by the fix).
  • Qwen3.6 MTP2 1k c=8 aggregate output 500 tok/s (440 before).

Final full test plan on f419d8d6c0 (2026-10-09 20:04 → 01:22)

15 runs, fresh cache each, TP=4 F&P unless noted, HIP 2-5/6-9. All gates PASS, 0 faults: probe before/after load, PPL, prefix, fault scan; concurrency probe PASS everywhere except Qwen3.6 PIECEWISE (prompt 6 token 17, margin 1.016) and DSV4 (solo-vs-solo repetition on the pruned checkpoint), both solo-vs-solo noise (the probe now classifies these as noise, 1c595eb646). Unit suite: 45 files pass; known failures only (listed in the checklist). Cells: prefill tok/s/req | decode tok/s/req | TTFT s | ITL median ms | aggregate output tok/s.

Run 1k/512 c=1 1k/512 c=8 16k/1k c=1 16k/1k c=8
27B AWQ 1099 · 44.0 · 0.93 · 22.7 · 40.8 256 · 23.1 · 4.01 · 36.6 · 156.9 894 · 38.9 · 18.3 · 25.7 · 23.0 404 · 6.5 · 40.5 · 52.8 · 41.2
27B AWQ W4A8=1 1436 · 44.1 · 0.71 · 22.7 · 41.6 337 · 24.1 · 3.04 · 36.5 · 168.7 1112 · 39.0 · 14.7 · 25.7 · 25.0 500 · 7.5 · 32.7 · 52.8 · 48.1
27B EXL3 MTP0 1232 · 34.3 · 0.83 · 29.2 · 32.5 287 · 21.1 · 3.57 · 41.4 · 147.3 1015 · 32.2 · 16.2 · 31.1 · 21.4 458 · 6.9 · 35.8 · 56.0 · 44.4
27B EXL3 MTP2 (greedy) 1241 · 81.0 · 0.82 · 36.2 · 71.8 390 · 30.3 · 2.63 · 69.9 · 202.0 953 · 65.1 · 17.2 · 45.6 · 31.1 420 · 7.8 · 39.0 · 102.1 · 48.2
Flash-Next MTP0 1511 · 43.4 · 0.68 · 23.0 · 41.1 497 · 25.1 · 2.06 · 36.1 · 182.3 1979 · 42.9 · 8.28 · 23.3 · 31.9 1033 · 12.9 · 15.9 · 36.9 · 85.7
Flash-Next MTP2 (greedy) ⚠ 1957 · 68.2 · 0.52 · 29.8 · 63.9 908 · 24.8 · 1.13 · 72.4 · 172.7 1948 · 79.3 · 8.41 · 30.2 · 48.0 1114 · 10.6 · 14.7 · 73.8 · 72.1
Flash-Next agentic (flashnext.env, SEQS=6) 1512 · 41.7 · 0.68 · 24.0 · 39.6 135 · 28.4 · 7.60 · 32.5 · 145.5 1971 · 41.2 · 8.31 · 24.3 · 30.9 451 · 15.7 · 36.3 · 33.2 · 76.4
Qwen3.6-35B-A3B MTP0 5041 · 75.2 · 0.20 · 13.3 · 73.2 1433 · 51.1 · 0.71 · 18.4 · 381.8 4002 · 66.4 · 4.09 · 15.1 · 52.5 1670 · 20.6 · 9.81 · 27.0 · 137.2
Qwen3.6-35B-A3B MTP2 (greedy) 4995 · 152.0 · 0.21 · 17.1 · 143.5 2113 · 59.1 · 0.48 · 40.9 · 439.9 3750 · 126.1 · 4.37 · 21.5 · 82.0 1686 · 23.4 · 9.72 · 58.0 · 151.3
Qwen3.6 PIECEWISE only 4987 · 33.0 · 0.21 · 30.4 · 32.6 1425 · 30.5 · 0.72 · 31.6 · 234.2 4002 · 32.8 · 4.09 · 30.4 · 29.0 1707 · 18.8 · 9.60 · 31.5 · 127.8
full.env 27B AWQ TP=2 559 · 29.9 · 1.83 · 33.5 · 27.0 61 · 20.2 · 16.8 · 45.0 · 89.2 491 · 25.5 · 33.4 · 39.3 · 13.9 122 · 4.8 · 134.6 · 79.8 · 22.7
Flash-Next eager 465 · 10.2 · 2.20 · 94.4 · 9.8 304 · 9.0 · 3.36 · 101.3 · 67.8 — —
27B AWQ eager / DSV4 gates only (DSV4 PPL 8.44; benches from earlier runs)

Per-request prefill at c=8 includes queueing behind the other prefills (aggregate prefill ≈ c=1 capacity, see bench_metrics.py); the c=8 decode/ITL-mean gap is the mixed-batch stall being fixed below. ⚠ Flash-Next MTP2 c=8 came in below the previous tip (1k c=8 201.0 → 172.7, 16k c=8 92.8 → 72.1 aggregate). A/B: the e7e63a0a75 gemv gate (M ≤ 24) costs ~9 % decode ITL at MTP2 c=8 (tuned rocBLAS beats gemv at M=16-24 on every Flash-Next shape) → reverted to M ≤ 8 in b074874b54 (kernel capacity + tests kept, gemv_crossover_bench.py added). The rest of the gap is run-to-run variance at c=8 (same tree: 172.7 vs 186.1 tok/s), i.e. the mixed-batch stall.

Cross-slot mamba/GDN state fix + backports (pushed: 1467ee7dc4..6fd30fd068)

Fix: our port of upstream vllm-project#55506 (4b8444c248) handed the aligned-state-index kernel per-slot block tables that it indexes by batch row, so when slot ≠ batch row a request read another slot's GDN/PLE state or the null block. Root cause of eager Flash-Next token salad + HSA fault (first eager request lands in slot 7 as batch row 0); in F&P the compile dummy runs masked it (residual exposure in the req_idx pre-copy, closed by the fix). With it, --enforce-eager works without the JIT-warmup workaround (d11df8c312). New standing gate: concurrency_probe.py (solo vs c=8 vs churn vs long prompts crossing a mamba block; classifies divergences by top-1/top-2 logprob margin + continuation quality; F&P Triton 0 fail / 8 near-tie, F&P HIP 0 fail / 17 near-tie, all within solo-vs-solo noise).

Run (TP=4, HIP 6-9) PPL Concurrency Cell Prefill tok/s Decode tok/s/req Agg out tok/s TTFT s ITL med ms
Flash-Next MTP0 F&P, Triton conv 4.86 PASS* 1k c=1 / 1k c=8 / 16k c=1 / 16k c=8 1574 / 516 / 2061 / 1054 44.1 / 25.4 / 43.5 / 13.1 41.8 / 184.8 / 32.6 / 87.4 0.65 / 1.98 / 7.95 / 15.55 22.7 / 35.9 / 23.0 / 36.5
Flash-Next MTP0 F&P, HIP conv 4.77 PASS* same 1549 / 502 / 2016 / 1050 44.1 / 25.2 / 43.5 / 13.1 41.8 / 183.2 / 32.4 / 86.8 0.66 / 2.04 / 8.13 / 15.61 22.7 / 36.0 / 23.0 / 36.6
Flash-Next MTP0 eager, Triton conv (no workaround) 4.86 PASS 1k c=1 / 1k c=8 1494 / 513 10.5 / 9.8 10.4 / 75.4 0.69 / 2.00 94.6 / 99.6
Flash-Next MTP0 eager, HIP conv (no workaround) 4.78 PASS 1k c=1 / 1k c=8 1480 / 513 10.8 / 10.0 10.7 / 76.7 0.69 / 2.00 92.4 / 97.0
Flash-Next MTP2 F&P 4.84 PASS 1k c=1 / 1k c=8 / 16k c=1 / 16k c=8 2031 / 883 / 2002 / 1230 65.3 / 28.0 / 82.6 / 13.9 61.4 / 201.0 / 49.8 / 92.8 0.50 / 1.16 / 8.19 / 13.32 29.7 / 67.0 / 30.1 / 67.6
27B AWQ F&P 4.94 PASS 1k c=1 / 1k c=8 / 16k c=1 / 16k c=8 1137 / 264 / 905 / 406 43.8 / 23.3 / 38.8 / 6.5 40.8 / 158.5 / 23.0 / 41.3 0.90 / 3.88 / 18.11 / 40.40 22.8 / 36.5 / 25.8 / 52.9
27B AWQ eager (no workaround; two-shot settings forced → clean fallback) 4.94 PASS 1k c=1 / 1k c=8 1113 / 284 17.2 / 12.8 16.7 / 93.9 0.92 / 3.61 58.2 / 68.7

* these two runs logged FAIL with the first, strict probe version (any divergence before token 24); re-run with the discriminating probe: 0 fail. All runs: probe + after-load probe + prefix PASS, 0 faults. HIP vs Triton conv update at parity → recipe default unchanged. Old-tree contrast: eager solo requests fail (salad, NaN logprobs → HTTP 400). Final run after the rebase over a43be93b10: F&P PASS, concurrency 0 fail / 9 near-tie. MTP=2 cold start F&P: decode warmup on/off identical (boot 927/923 s, first TTFT 3.18/3.17 s). 145 unit tests pass.

Fixed during validation (bugs from upstream API changes that the merge didn't surface):

Commit Fix
8202cb647b RDNA2 W4A16 kernels ported to the post-vllm-project#54809 GPTQ API (no act-order); scheduler kv_holding_waiting; MXFP4 output_topk
0f592bea03 Qwen4Exp PLE accepts compressed-tensors checkpoints (Flash-Next W4A16 failed to boot)
e5ecbfd20a RDNA2 W4A16 MoE: restore K-first weights after vllm-project#52798's N-first CT buffers (profile run died with [2048, 80] vs [2048, 2560])
3f3776778b Hybrid W4A16 gfx10 dense dequant reads vllm-project#54965's packed 4-bit zero-points (was all-NaN for asymmetric layers)
30f72c2291 BF16x3 router warmup gated on CUDA (gfx10 capability (10,x) aliases SM100 → No module named 'cutlass')
bc882a3a45 Resident MoE path skips the new N-first transpose (it was transposed twice)
523eb846e0 MXFP4 dense + MoE kernels (DSV4 MXF4_RDNA2): accept 32-bit [..,K/8,N] or 8-bit [..,K/8,4N] weights and uint8 or float8_e8m0fnu scales. Also fixed three dequant bugs hidden behind the old dtype check: LUT returned ints as half (0.5 → 14336.0), scale row stride ×4, wrong per-column scale lanes. Predates the port (on 0.28 every call was rejected)
cd8056fd51 Qwen4Exp PLE on RDNA runs as a graph-splitting op. Upstream's in-graph n-gram ids were corrupted on padded piecewise HIP-graph replays: every in-sequence row of a prefill padded up to a capture size diverged, and MTP2 answered 1+1=3. Localized with per-row GDN traces
6e56ecb782 GDN update_block_table on RDNA kept upstream's paged-slot semantics while build() used state arenas. Extra GDN KV groups got paged slot ids as arena rows plus the source group's cache slots. Arena rows now land in the group's own buffer and only its paged slots are re-gathered
0550c0f655 Removed the stale qwen4_exp/amd/hc.py duplicate (double op registration)
66cb4ab3a9 27b-awq recipe VLLM_USE_AOT_COMPILE=0 like the other recipes. AOT hands rank 0's graph to all ranks on the qwen3_next path; same failure on 0.28
d11df8c312 --enforce-eager on RDNA keeps the JIT warmup. Without it, Flash-Next returned NaN logits. Superseded: root cause was the cross-slot state bug (5bf02f78ea), workaround reverted in eb503cd3ca
f41482de7b RDNA2 W4A16 MoE for GPTQ (MoeWNA16) checkpoints. Failed on 0.28 too ("Unsupported wna16 MoE backend"), and Triton WNA16 gives garbage on gfx1030. Kernel matches an fp32 dequant of real experts (0.0015 / 0.0002)
1578e258c9 QSA rope/rmsnorm fast paths gated on RDNA (were on for any ROCm) and limited to device tensors
bda1589fc6 Mamba block-aligned splits for all V2-runner hybrids at TP>2 → prefix hits with MTP on 27B/EXL3 (were 0 %)
00fcc97ff5 rdna_ar one-shot all-reduce returned one shared persistent output buffer, so every all-reduce clobbered earlier results still in use (the embedding all-reduce output kept as the layer-0 residual was overwritten by the first token-mixer all-reduce). Hits every TP>1 model when the message fits the 64 KiB one-shot gate (short prompts, decode). 0.28 has the same bug. Fresh per-call output (empty_like, kept alive under capture); new tests/distributed/test_rdna_ar_gpu.py. Root cause of the Qwen3.6 garbage
3e9ca702a1 RDNA kernel layout: csrc/rocm/rdna/common/{arch,convert}.cuh (compile-time RDNA2/3/3.5/4 traits, bit-exact fp16/bf16/fp8/e8m0 conversions); DSV4 sparse-MLA + indexer HIP kernels moved to csrc/rocm/rdna/dsv4/ (op names unchanged); vllm/models/deepseek_v4/amd/rdna/ (wrappers + fakes)
01b439bfd2 DSV4 q-norm + RoPE + KV-insert RDNA HIP kernel _rocm_C::dsv4_qnorm_rope_kv_insert_rdna: fp16 or bf16 q/kv in, upstream fp8 DS-MLA cache row out (448 fp8 + UE8M0 scales, 64 bf16 RoPE). 93/93 tests: cache bytes identical to upstream _C for bf16 inputs, cudagraph replay = eager, rows decode correctly through sparse_mla_decode_rdna2
8b06c28dfb RDNA2 GEMM/attention ops returned one shared persistent output buffer: gptq_gemm_rdna2 / gptq_gemm_rdna2_prefill (every W4A16 dense layer) and the fp8/int8 FA-RDNA2 entries zeroed and overwrote the previous call's output, so two outputs alive at once, or an output fed back as the next GEMM's input after an in-place norm (eager), were clobbered (new test reproduces both on the old build). Fresh per-call outputs; split-K scratch stays persistent. Also restores upstream allocation in LLMM1 / wvSplitKrc / wvSplitK_int4_g (they run on gfx9/gfx11 too). Serve A/B on F&P: Flash-Next MTP0 and 27B AWQ probe/prefix/after-load PASS, all 8 cells within noise (Flash-Next 16k c=8 84.7 vs 81.6, 27B 16k c=8 44.1 vs 44.0 tok/s)
ea8785b323, b08d5dc7bd ruff format / clang-format of fork-owned files (hipify fix for one split kernel launch)
2cc2fbc7c6 HIP gated RMSNorm skipped on every GDN layer: upstream now passes core_attn_out/z as 3D [tokens, heads, head_dim]; the RDNA forward_hip only took 2D, so all 48 GDN layers fell back to the decomposed native norm (~9 kernels each). Flatten/restore around the HIP kernel. 27B AWQ decode model forward 22.64 → 20.66 ms/step (0.28: 21.50). Root cause of the 27B c=1 gap vs 0.28; v0.31 now beats 0.28 on all 4 cells (agg out 1k c=1 40.7 vs 39.7, 1k c=8 157.3 vs 157.2, 16k c=1 22.9 vs 22.6, 16k c=8 45.2 vs 41.3). New test in tests/kernels/mamba/test_gdn_output_projection.py
b7384d1aca serve_validate.sh BENCH_EXTRA (e.g. --temperature 0). The EXL3 MTP=2 "regression" was a bench artifact: the 0.28 reference used greedy; benched the same way (greedy, 4 cells) v0.31 is at parity (1k c=8 44.5 vs 44.5, 16k c=8 29.7 vs 28.4 tok/s, acceptance 2.84 vs 2.79)
7b98e3f851 MXFP4 E2M1 lookup table was wrong for ±3/±4/±6 (fp16 bit patterns of 4/8/16 instead of 3/4/6, introduced with the dtype fix in 523eb846e0; its test only drew weights in [-1, 1]). DSV4 MoE outputs ~1.7× too large, compounding to fp16 NaN by layer ~34: root cause of DSV4's garbage output. Test now covers [-6, 6] and checks the table; relative error vs fp32 dequant of real checkpoint experts 0.8 → 1e-3
5eb773b656 RDNA2 MXFP4 MoE applies swiglu_limit (clamp like SiluAndMulWithClamp)
6f5333e868 sparse_mla_prefill_rdna2 staged 2048 float4 per KV row instead of 64 (out of bounds in global and shared memory)
7df02835a2, 6093e3a899 DSV4 fp16 attention, plan steps 3-5: inverse RoPE / wo_a / indexer q-quant follow the activation dtype; sparse indexer on gfx10x with HIP paged MQA logits (C4A block-flat layout, graph-safe split grid); no 4 GiB profile workspace; fp16 decode wrapper; DeepseekV4RDNAAttention selected on RDNA (CDNA unchanged). test_attention_ops.py 13/13, test_kv_insert.py 93/93, test_mxfp4_rdna2.py 51/51
08adedcc95 ppl_probe.py: perplexity on a fixed paragraph + token-salad check, run by serve_validate.sh (PPL_PROBE, MAX_PPL), for degraded/pruned checkpoints where exact-answer probes are not meaningful
36c4164b22 Block-aligned mamba splits only for MTP and Flash-Next on V2 (bda1589fc6 had enabled them for every V2 hybrid). With EXL3's 784/800-token state blocks a 1k prompt prefilled in 2 steps and 16k chunks shrank 2048 → 1568: EXL3 MTP0 1k c=1 median TTFT 1148 → 824 ms, 16k c=1 prefill 961 → 1029 tok/s (0.28: 1005; this was the "EXL3 16k prefill 7 % behind 0.28"). MTP0 prefix hits still work via the prefill checkpoint (31.5×). 27B AWQ / Flash-Next (1024-token blocks) unaffected
9504d8091c RDNA2 fp8 E4M3 decode turned 256..448 into NaN (treated exponent 0xF as Inf/NaN; e4m3fn has no Inf, only S.1111.111 is NaN). Shared by the W8A16 fp8 dense/MoE, W8A8 fp8 dense and DSV4 indexer kernels
eec301a25a RDNA2W8A16Fp8BlockLinearKernel registered and selected for fp8 block linears on gfx1030 (v0.31 never registered it: every fp8 linear ran Triton, which emulates fp8 MMA there, 1.3-3.9 ms vs ~0.2 ms per decode GEMM); weight/scale prep once at load. DSV4 decode 1.3 → 4.2 tok/s
dcd2388100, 14678c9e0e DSV4 indexer prefill MQA logits memory-bounded (torch fallback built a 64× oversized scores tensor, OOM at 16k); dsv4-flash.env on FULL_AND_PIECEWISE with breakable graphs
ec349e9079..c19e8c0c15 DSV4 decode 4.2 → 14.6 tok/s/req, prefill 142 → 960 tok/s (1k c=1, F&P): fused HIP mHC pre/post (torch fallback was ~130 kernels/call, 172 calls/step); MXFP4 MoE block size from the token count; gemm_w8a16_fp8_dense 7-38× (dword loads, branch-free fp8, dequant once per M tile); branch-free E2M1 in the MXFP4 GEMMs (10-14×); compressor score GEMMs on gemv_f16_rdna2; split-KV sparse_mla_decode_rdna2 (1.35 → 0.061 ms). PPL 8.30-8.49 throughout
0cfce962e9 Legacy recipes on the validated config: flashnext.env (agentic: tools/reasoning/vision, 0.0.0.0, SEQS=6) V1 runner + breakable + 16-token blocks → V2 + F&P + RDNA_ATTN + 1024-token blocks + mamba-align (probe/PPL 4.82/prefix PASS, 1k c=1 39.9 tok/s); full.env default GPUs 0,1 (W7800s) → 2,3 (27B AWQ TP=2 PASS)
5148ba7817..59d8d6b067 DSV4 decode 14.6 → 29.7 tok/s/req (1k c=1), 1k c=8 aggregate 48 → 129 tok/s: branch-free exact E4M3 decode in qdq_fp8_rdna2.cuh (+ full-range E4M3 test); wo_a per group on gemm_w8a16_fp8_dense from the fp8 checkpoint weight (was a dequantized fp16 einsum; −0.7 GB); fused clamped SwiGLU in the MXFP4 MoE (bit-identical); C128 compressor decode on the two-stage split (155 → 11 µs/launch, byte-identical rows); recipe SEQS=8 / KV=6e9 (c=8 no longer queues; 8×17k peaks at 40 % KV). PPL 8.46, probes + prefix PASS. Prefill ~4 % lower than round 1 (wo_a path; 960 → 931 tok/s at 1k)
c2cbf00477, 0835e3cbf5, d1c3e6c5d3, 78d62cf7c1 27B EXL3 MTP=2 c=8: 44.5 → 207.6 tok/s aggregate, ITL 465 → 67 ms; c=1 decode 25.9 → 86 tok/s/req (PPL unchanged 4.95). Root causes: above M=8 the EXL3 GEMM fell back to a generic kernel that decoded each trellis tile 3× (332 ms of GEMM per step at M=24 → 29 ms with the new decode-once gemm_exl3_v3_kernel_rdna, gfx103x, bits 3/4, M ≤ 64, opt-out VLLM_EXL3_GEMM_V3=0); the 6 bpw lm_head also ran the generic kernel at every M (18-30 ms/call, 3 calls per MTP step → 0.6-1.0 ms). Recipe 27b-exl3 on rdna_ar one-shot; VLLM_RDNA_PHASE_TIMING debug timer. MTP0 also faster: 1k c=8 87.6 → 136.4, 16k c=8 36.9 → 43.3 tok/s. test_exl3_rdna2 533 passed. (the held gemv_f16_rdna2 M ≤ 24 commit landed as e7e63a0a75 after its check)
e7e63a0a75, f419d8d6c0 gemv_f16_rdna2 for M ≤ 24 on gfx10x (unquantized fp16 linears at MTP-verify sizes; lm_head at 24 rows 5.5 → 3.6 ms; no regression on Flash-Next / 27B AWQ MTP0: Flash-Next agg 42.1/186.7/32.7/87.5 tok/s) + EXL3 v3 GEMM tile-word prefetch (3-12 %). 27B EXL3 MTP=2 final: 16k c=1 decode 23.9 → 64.3 tok/s/req, ITL 124 → 46 ms; 16k c=8 ITL 497 → 102 ms, agg 29.5 → 53.0 tok/s; 1k c=8 agg 45 → 200 tok/s; concurrency 0 fail. 587 unit tests pass
b074874b54 gemv_f16_rdna2 dispatch back to M ≤ 8 on gfx10x: at M = 16-24 tuned rocBLAS wins on every Flash-Next shape (e.g. 10240×2560 at M=24: 177 vs 245 µs); Flash-Next MTP2 c=8 ITL 72.9 → 66.5 ms
a25f2e0b28 TunableOp coverage campaign into the single rocm7.14-rocblas5.5 profile: census (record_untuned, probes + 4 cells) over Qwen3.6 MTP0/MTP2, DSV4, 27B EXL3 MTP2, 27B AWQ W4A8 → 846 novel GEMM keys → tuned + curated (adopt ≥3 % over current and heuristic) → 622 new rows (979 → 1601/rank); per-shape median 1.79× (p90 3.02×, best 15×); lookup-hit 1601/1601; profiles.json covers all serving configs. End-to-end gain is small (Qwen3.6 1k c=8 agg 381.8 → 384.9 tok/s, TTFT 0.71 → 0.67 s). 6 shipped Flash-Next rows measured 3-10 % slower than the heuristic are kept and listed in provenance.json for review
2907c75d40..092d7ea70c DSV4 round 3: 1k decode 29.7 → 38.8 tok/s/req (ITL 25.8 ms), 16k prefill 732 → 1092 tok/s (TTFT 22.4 → 15.0 s), 16k c=8 aggregate 36.2 → 55.0 tok/s: W8A16 fp8 linears use per-call dequant + rocBLAS above M=256 (no persistent copy); mHC pre decode split-K + 16-lane Sinkhorn; tiled v_dot2 indexer decode logits (220 → 15.6 µs); radix top_k_per_row_decode on gfx10x (16× at long context); fused tiled indexer prefill logits (_rocm_C::mqa_logits_prefill_rdna2). Edge-case tests (context 0/1/31/33/4097/16383, -1 padding, masked tiles). Soak: 5 fresh boots incl. cold, 480/480 requests, 0 hangs/faults (RCCL); PPL 8.31
a43be93b10 rdna_ar one-shot wedge fixed: each wait spun a fixed 2e6 polls (~0.8 s on gfx1030; the "~1953 ms" in the abort was a poll count), so any rank reaching a collective late (first-request Triton JIT, long host step) made the peers abort, kill the engine and leave the marker that forces RCCL on later boots. Waits are now bounded by wall clock (VLLM_RDNA_AR_WAIT_MS, default 120 s, below the engine timeout); completed late peers are logged once (new op rdna_ar_slow_info), never fatal; the abort record holds real time. Reproduced with a 5 s injected skew (old: WEDGED, engine dead; new: completes, 32/32). 10 fresh boots incl. 2 cold, 0 wedges; unit tests 31 passed; DSV4 4 cells at parity
b03bf9d210 W4A8 validated on 27B AWQ (see the checklist); W4A8 fused-entry unit test no longer assumes a deterministic W4A16 split-K arm
5bf02f78ea Cross-slot mamba/GDN state bug (port of vllm-project#55506 bound per-slot block tables to a kernel indexing by batch row): eager first request read the null block → token salad + HSA fault; latent F&P exposure in the cross-block pre-copy. Fixed; 2 new GPU tests
eb503cd3ca Revert of the eager JIT-warmup workaround (d11df8c312): --enforce-eager now correct on Flash-Next (both conv paths) and 27B without it
1467ee7dc4 GDN causal_conv1d_update_rdna2 reads its state through all strides (backport of v0.29 271fae375b); 28 new tests; HIP vs Triton conv update at parity, recipe default unchanged
f7d4fa738c GDN decode / MTP-verify Triton kernels warmed at profile time on RDNA (backport of f923976302; no cost/gain on F&P, covers paths without the V2 JIT warmup)
611063c62a Broken two-shot all-reduce removed from rdna_ar (one-shot ≤ 64 KiB, RCCL above; ALGO=twoshot / large MAX_KB now warn once and fall back)
d95428e893 concurrency_probe.py: greedy solo vs c=8 vs churn vs long prompts, divergences classified by logprob margin + continuation quality; on by default in serve_validate.sh
— Driver env in spawned EngineCore/workers checked: nothing lost on v0.31 (backport item dropped)
— --enforce-eager correct on 27B AWQ with the fresh-output fix; the Flash-Next eager dependency on the JIT warmup was later traced to the cross-slot state bug (5bf02f78ea) and the workaround reverted (eb503cd3ca)
b36211bc19, a9c9fa106c, c02e5c96c8, 66291b8bc2 Fork tests: vllm-project#54809 kwargs, _get_linear_backend(quantization=), MoE W4A16 fp32 reference (was the dense op on MoE-layout weights), MXFP4 references + 32/8-bit/e8m0 coverage

Results

  • Build + registration: 505 s, 73 _rocm_C ops (58 RDNA), no unresolved libs

  • Flash-Next AWQ W4A16, TP=4 (HIP 2-5, PLX-B), F&P, RDNA_ATTN, upstream PLE, RDNA2 W4A16 MoE, rdna_ar one-shot active, 0 faults. Numbers are with the PLE fix:

    Cell v0.31 MTP0 0.28 MTP0 ref v0.31 MTP2 0.28 MTP2 (same day)
    1k/512 c=1 38.63 38.38 61.11 51.38
    1k/512 c=8 174.97 166.76 189.37 —
    16k/1k c=8 78.08 (TTFT 18.1 s) 74.2 (TTFT 41 s) 80.52 74.41
    prefix warm TTFT (12k words) 0.38 s (18×) — 0.94 s (8.9×) 3.76 s (2.9×)
    greedy probe before/after load 3/3, 3/3 — 3/3, 3/3 3/3

  • First-token logprobs on padded prompts now match the no-cudagraph run (e.g. 7 tokens: −0.715 vs −0.695). A small residual remains at 5 tokens (−0.78 vs −0.66): open

  • Unit, re-run after the fixes: test_rdna2_w4a16 18/18, _awq 8/8, _selection 18/18, test_rdna2_w4a8 18/18, test_rdna_hybrid_w4a16 76 passed, test_rdna2_moe_w4a16 75/75 (36 failed on 0.28), test_moe_wna16 25/25, test_mxfp4_rdna2 50/50 (29/29 failed on 0.28), moe_w4a8 27/27, exl3 215/215, rocm_moe_skinny 57/57, GDN decode/prep pass

  • Qwen3.8-27B AWQ dense, TP=4, F&P, AOT off (needed on 0.28 too): probe 3/3 before and after load, prefix hit 25×, 0 faults, RDNA2W4A16LinearKernel + RDNA_ATTN

    Cell v0.31 0.28 (same day)
    1k/512 c=1 36.68 39.27
    1k/512 c=8 149.08 154.46
    16k/1k c=8 40.37 40.86

    ✅ Resolved: the gap was the HIP gated RMSNorm being skipped on every GDN layer (2cc2fbc7c6); v0.31 now beats 0.28 on all 4 cells. Re-checked with aligned splits on: 36.58 / 149.54 / 40.39, probe 3/3 before and after load, prefix 24.7×, 0 faults

  • test_gdn_prefill_solve_wy_rdna2 22/22, test_gdn_metadata_builder 53/53 (5 failed before the fix), test_gdn_prefill_rdna2 7/7, FA-RDNA2 shape sweep 16/16 + writer layout 56/56, test_rdna_p2p 19/19, test_kv_cache_utils 167/167, test_prefix_caching 169/169, test_prefill_cadence 6/6, kv_offload/test_factory 48/48, test_qwen4_exp (spec) 2/2

  • Still failing, also on 0.28: test_rocm_skinny_gemms aborts; test_gdn_prefill_kkt_rdna2 fails and then hangs the GPU

  • Also failing on 0.28 (not caused by the port): test_gdn_prefill_delta_h_rdna2 29/30, test_gdn_prefill_o_rdna2 30/42 (both the opt-in HIP GDN prefill chain, off by default); test_fused_shared_expert 12 (MXFP4 linear kernel unavailable); compile/test_config (tests that assume a CUDA host; more of them upstream now); test_partial_prefix_cache_hits 5 (stale scheduler mocks)

  • Expected after the PLE decision: test_ple_offload_worker, test_ple_quant_gather (sidecar unwired). To be skipped on v0.31

Cudagraph policy: eager-only paths count as bugs. Every model must run correctly on FULL_AND_PIECEWISE.

  • Recipes: all F&P by default, validated (Flash-Next MTP0/MTP2, 27B AWQ, 27B EXL3 now F&P instead of eager, full). DSV4 boots and serves on F&P (see the DSV4 row).
  • No runnable model needs --enforce-eager any more.
  • Eager islands inside F&P (@eager_break_during_capture) are inactive in all validated recipes: our checkpoints resolve to *ForConditionalGeneration, which isn't on the breakable-arch list, and the recipes don't set VLLM_USE_BREAKABLE_CUDAGRAPH, so GDN/QSA/FA ops were fully captured in every F&P run. flashnext.env is now on V2 + F&P like the others (0cfce962e9); full.env still forces breakable graphs on. Islands, for reference:
    • [backlog] GDN full-layer op _rdna_full_forward + Qwen3.5/3.8 breakable-arch opt-in (original reason: GDN prefill chain faulted under capture; re-test on v0.31)
    • [backlog] FA-RDNA2 backend op (rdna_attn.py)
    • [backlog] QSA op (qwen4_exp/amd/qsa.py)
    • [backlog] (ple_offload_layer.py: sidecar, unwired; upstream sparse-MLA indexer: DSV4)

Open / suspected upstream issues

  • Upstream _process_weights_rdna3 (gfx1100) still assumes [E, K//8, N] after [Quant] Use canonical N-first weight format for CT WNA16 MoE vllm-project/vllm#52798; probably the same bug as our RDNA2 fix (needs a gfx1100 check)

  • ROCm gfx10 aliases capability family 100 in other checks too (already in 0.28, e.g. Triton unified attention picks B200 tiles for head_size 256)

  • triton_kernels import fails on triton 3.7.1 (is_hip_gfx1250); only the upstream Triton MXFP4 MoE path is affected

  • One host crash during the first Flash-Next boot, with PLX-B AER bursts while unit tests ran at the same time; not reproduced with isolated runs

  • Dense gptq_gemm_rdna2 fed MoE-layout weights is ~3 % off fp32: not a valid MoE reference. The MoE test now uses an fp32 MoE reference (b36211bc19..) and the dense layers pass their own tests

  • Qwen3.8-27B EXL3 3.0 bpw (mul1), TP=4, FULL_AND_PIECEWISE cudagraphs (now the recipe default): probe 3/3 before and after load, prefix hit 31×, 0 faults, boot 181 s (eager: 701 s). Eager also passes 3/3

    Cell v0.31 F&P 0.28 F&P (2026-10-01)
    1k/512 c=1 18.52 18.91
    1k/512 c=8 84.06 81.77
    16k/1k c=8 35.99 30.58

  • EXL3 27B MTP=2 on F&P cudagraphs: probe 3/3 before and after load, mean acceptance length 2.6, 0 faults

    Cell v0.31 MTP=2 0.28 MTP=2 F&P
    1k/512 c=1 21.25 22.05
    1k/512 c=8 36.73 45.28
    16k/1k c=8 24.21 24.52
    • ✅ c=8 gap resolved: a bench artifact (0.28 reference was greedy; at parity benched the same way, b7384d1aca), then 44.5 → 207.6 tok/s aggregate with the EXL3 v3 GEMM (c2cbf00477..)
    • ✅ Prefix caching with MTP=2 fixed (bda1589fc6): mamba block-aligned splits are now on for all V2-runner hybrids at TP>2 (the RDNA workaround kept them for Flash-Next only, unchanged since 0.28). Result: 6.8× hit, identical cold/warm text, probe passing before and after load, 0 faults; 1k c=1 21.2, c=8 35.9, 16k c=8 23.5 tok/s

  • DeepSeek-V4-Flash K128 (MXFP4 experts, ~50 % expert-pruned checkpoint): serves coherently in fp16 at TP=4 on FULL_AND_PIECEWISE (eec301a25a): perplexity 8.54 (eager 8.4), probe 3/3 before and after load, prefix PASS (16k cold 115 s → warm 2.0 s), rdna_ar on, 0 faults. It never ran on 0.28. Perf (092d7ea70c): 1k c=1 decode 38.8 tok/s/req, ITL 25.8 ms, 16k prefill 1092 tok/s (first working run: 4.2 tok/s, 238 ms, 142 tok/s prefill); 4 cells in the table above. Earlier fixes:

  • Qwen3.6-35B-A3B GPTQ-Int4, TP=4, F&P, RDNA2 GPTQ MoE kernel + rdna_ar fix: probe 3/3 with coherent long replies, prefix hit 12.2×, after-load probe passes, 0 faults

    Cell Prefill tok/s Decode tok/s/req Agg out tok/s TTFT s ITL med ms
    1k/512 c=1 4627 69.5 67.6 0.22 14.4
    1k/512 c=8 1219 48.2 357.7 0.84 19.4
    16k/1k c=1 3911 61.7 49.3 4.19 16.2
    16k/1k c=8 1788 20.0 135.1 9.16 27.9

Done. Next (separate work): Docker image blivioniag/vllm-rdna:v0.31.0-extras from rdna_extras (opengfx1030/vllm-rdna-docker, same rocm-rdna:7.14.0 base, no patches).

Mixed-batch stalls

Reproducible probe landed (5890e7fea0..7d9244387a): tools/rdna/port_v031/stall_probe.py (decoders in steady state + timed long-prefill injections, reverse case, repeats, gates; STALL_PROBE=1 in serve_validate.sh; stall_ab.sh for knob A/B), doc docs/rdna2/mixed-batch-stall-probe.md.

Baselines, 16k injection, 6 decoders, TP=4 F&P (decoder ITL baseline 34-35 ms):

arm max gap gap p50 frozen decode tok/decoder during prefill prefill TTFT × solo reverse: decoders blocked
Flash-Next v0.31 default 515 ms 508 ms 0.99 19 1.11 7.1 s
Flash-Next 0.28 default 547 ms 522 ms 0.99 18 1.10 7.0 s
Flash-Next interval 4 + LPTH 256 247 ms 34 ms 0.70 260 3.01 2.7 s
27B AWQ v0.31 default 2672 ms 1978 ms 0.99 11 1.03 17.8 s

Stall fix round 1 (pushed: a5a456939a..8aa453030d): per-op GPU/host timing of mixed steps showed (a) every step > 8 tokens runs with no cudagraph (capture sizes stop at 8): host launch ≈ 235 ms flat, (b) GPU ≈ 95 ms + 0.40 ms/token across MoE, GDN, hyper-connection mix, all-reduce and fp16 GEMMs; host-bound below ~450 tokens, GPU-bound above. Landed: GDN state arenas sized by the decode batch instead of the capture ladder (was 1.5 GiB/layer, blocked larger captures); hyper-connection prefill GEMM rows padded to multiples of 64 + 196 new TunableOp rows in rocm7.14-rocblas5.5 (783 → 979/rank; HC down GEMM at 512 rows 0.89 → 0.23 ms); host/op/sync debug timers (off by default). Ruled out on GPU: host syncs, mamba pre-copy, PLE, autotune (JIT only at first admission of a new QSA width, plausibly the 20.6 s outlier); 16/32-row MoE tiles were 1.4-5× slower (reverted). Re-measured after round 2 and landed for dense models (decision 2026-10-10): a scheduler cap bounding prefill tokens in steps that carry decodes + decode-first ordering (27B max gap 2.67 s → 0.39 s, reverse 17.8 → 4.8 s; Flash-Next reverse 7.1 → 3.2 s, inject unchanged; but 27B 16k/1k c=8 aggregate −9 %). Next: graphs for mixed sizes (needs the RDNA graph wrapper to share activation buffers across capture sizes; the 2048 ladder OOMs today), fewer launches in the opaque GDN layer, non-replicated hyper-connection prefill.

Stall fix round 2 (pushed: dc9e536f13, 0095c9d00b, defaults unchanged): cudagraphs for mixed/prefill sizes via piecewise-graph input buffers shared across capture sizes (opt-in ladder via CG_SIZES, e.g. up to 2048 incl. 1032; capture 56 s / 4.2 GiB): host time of a 1024-token step 235 → ~98 ms, small mixed steps 169-207 → 74-88 ms, correctness PASS; and an opt-in VLLM_RDNA_GDN_OPAQUE=0 that keeps the GDN projections in the compiled graph (2048-token step host 94 → 80 ms). But Flash-Next mixed steps at the default 1024-token chunk are GPU-bound (~95 ms fixed + 0.40 ms/token: MoE 0.10, GDN 0.08, hyper-connection 0.07, all-reduce 0.08, QSA 0.06 ms/token), so the stall gates don't move (max gap ~0.53 s). Not for 27B: the ladder costs 13 GiB of graph memory and c=1 ITL 23.4 → 26.2 ms (replay-copy wrapper, being removed). Hyper-connection token-split across ranks: net ~3 % only (needs an all-gather), not done; fusing away its gate tensor ≈ 5-10 ms/step. Scheduler cap re-measured on this tip (27B: max gap 2.67 s → 0.33 s, decoders 8× more tokens, reverse 17.9 → 4.8 s, 1k c=8 TTFT 4.0 → 3.0 s; cost 16k/1k c=8 TTFT +14 %, decode −2 %; Flash-Next neutral without the ladder, reverse 6.9 → 3.2 s). Landed for the dense models only (38d2a446ec, 5d9f76f144): off by default (VLLM_RDNA_DECODE_STALL_MS=0), set to 250 ms in 27b-awq, 27b-exl3 and full.

Round 3 (merged 9cc6d6c61a..3b43339ab0; validated in the final pass):

  • Fused hyper-connection prefill — merged (9cc6d6c61a..946668b3d3; MTP=2 serve check still to run as a confirmation): new kernels fuse silu + up GEMM + sigmoid + gated mean (the [M,4H] gate never hits memory) after a split-K HIP down GEMM; on by default on gfx1030 in that branch (VLLM_RDNA_HC_PREFILL_FUSED=0 opts out). Per HC mix 1.47 → 1.32 ms at 2048 rows, 2.0× at 16 rows. The old 0.28 HIP prefill path was ~8× slower than torch (no fusion, every thread recomputed the whole row): kept opt-in and documented. Flash-Next MTP0: 1k c=1 prefill 1913 → 2041 tok/s/req (TTFT 0.54 → 0.50 s); 16k c=1 2051 → 2151 (TTFT 7.99 → 7.62 s); 16k c=8 aggregate prefill 3705 → 3864, aggregate output 86.8 → 88.7; decode and ITL unchanged. All gates PASS, 34 unit tests. Down projection not bit-reproducible run to run (fp32 atomics). Default on for gfx1030; other arches and non-gfx1030 builds keep the torch path.
  • Decode-stall cap fit (Flash-Next opt-in) — merged (f8af9e1b89, 9239abcf83; settings and trade-off tables in mixed-batch-stall-probe.md): the step log showed the earlier arms still ran 1024-token mixed chunks. The cause was not the mamba block rounding (512-token chunks work) but the cap's fit: prefill-only steps pushed the fitted cost above the budget. Fix: only steps with decoders feed the fit. Measured (budget 250 ms, efficiency floor 1.0, ladder to 1032): inject-16k max decode gap 505 → 341 ms, decoder tokens during the prefill 115 → 215, prefill 1.52× solo; cost 16k c=8 aggregate output 85.1 → 74.1 tok/s, TTFT 15.9 → 19.2 s. The ≤300 ms gate still fails because the cap could not drop below its initial 512; second fix (one bucket may shrink the cap, expected ~384-token chunks, ~280 ms steps) unit-tested only. Also a warm-up over all capture sizes before measuring (first use of a size compiled mid-bench). On in the dense 27B recipes (2672 → 394 ms max gap for −9 % 16k c=8 output), opt-in for Flash-Next (commented in its recipes) because the trade-off is worse there. The fit change also applies to the dense recipes; the final pass re-measures 27B with it.
  • Replay-copy wrapper removed — merged (ab72778560..3491035be0: GDN opaque-forward output allocated in the graph, which the wrapper was hiding (27B died at warmup with any ladder above 8); all-reduce bench; wrapper removal, all-reduce back inside the graph pieces as upstream). Graph memory 3-6× lower at parity speed (Flash-Next 0.86 → 0.13 GiB, 27B 2.8 → 0.44 GiB at default sizes). with the replay copies off, Flash-Next MTP0 (default capture sizes) and 27B AWQ (ladder up to 512) pass every gate incl. concurrency; 27B ladder graph memory 13 GiB → 1.44 GiB and c=1 ITL median back to 22.9 ms (26.2 with the wrapper). Open: Flash-Next MTP0 with the 512 ladder had one concurrency FAIL (c=8, prompt 1 token 25, margin 2.94, fluent continuation, not seen solo-vs-solo); one repeat on a fresh boot PASSED (the same token is now a soft warning, margin 0.27; greedy and after-load probes pass), so it did not reproduce in 1/1 repeats. The final pass re-runs the 512-ladder concurrency probe twice, Flash-Next MTP=2 and the default sizes on the merged code.
  • All-reduce (microbenchmarks only, no serve A/B yet; docs/rdna2/tp4-allreduce-sizes-2026-10-10.md): at prefill sizes RCCL is already at the PCIe link limit (23-25.6 GB/s bus bandwidth; 8 MB 501 µs, 20 MB 1228 µs), so no RCCL setting or custom kernel can beat it (the 2026-10-06 "0.06 ms at 10-21 MB" was an unsynced timing). LL128 = default, Tree worse, LL only wins below ~100 KiB. The earlier "+5 % with 8 channels" was noise: RCCL ignores NCCL_MIN_NCHANNELS below 8 GPUs, so it is not landed. Lever found: the one-shot rdna_ar beats RCCL up to ~240 KiB (80 KiB: 45 vs 65 µs in a graph; 96 KiB: 49 vs 69 µs), while the gate is 64 KiB, so 27B c=8 decode (80 KiB) and Flash-Next MTP2 c=8 verify (96 KiB) run on RCCL. Raising the gate to 128-160 KiB should save ~15-20 µs per all-reduce (~2-3 ms per 27B c=8 step); needs a serve A/B before landing. A reduce-scatter + allgather kernel could only win at 256 KiB-2 MiB (~1 % of a mixed step): stays in the backlog. TP=4 ring is 47→4A→52→55, all P2P, one PLX switch tree per set; no new PCI SERR/PERR during the sweeps.

Same behaviour on 0.28 (not a port regression). Lead: any step carrying prefill pays ~200 ms on top of the 33 ms decode step regardless of chunk size (256 tokens ~240 ms, 1024 ~475 ms); in mixed steps mamba-align rounds the leftover budget down to a 1024-token chunk. Suggested gates for fix rounds: max stall ≤ 300 ms, ≥ 100 decode tokens/decoder during a 16k prefill, prefill TTFT ≤ 1.5× solo, reverse stall ≤ 3 s.

Found along the way, fixed (a43be93b10): the rdna_ar one-shot all-reduce wedged when one rank reached a collective late (its wait cap was ~0.8 s). Now wall-clock bounded (120 s) and late peers are non-fatal.

Backports from the older forks (research 2026-10-09)

Done (1467ee7dc4..6fd30fd068): GDN conv-update stride fix · GDN decode/MTP warmup · rdna_ar two-shot removed · driver env checked (nothing lost, dropped) · plus the cross-slot mamba-state fix and the eager-workaround revert found along the way.

Future PR: rdna_ars reduce-scatter + allgather all-reduce (TP=8 / prefill payloads): code saved on wip/rdna-ars-snapshot-2026-10-06, plan in docs/rdna2/rdna-ars-allreduce-plan.md.

Secondary backlog:

  • Agentic flashnext.env on the mtp0 performance block: done in 20edef5321 (first try hit a one-off host LLVM ERROR: out of memory; bisect passed 3/3, full validation passed)
  • Raise the one-shot rdna_ar gate 64 → 128-160 KiB (microbenchmark: beats RCCL up to ~240 KiB; 27B c=8 decode and Flash-Next MTP2 verify sit at 80-96 KiB); needs a serve A/B
  • Flash-Next stall ≤ 300 ms: GPU-check the single-bucket cap fix (expected ~384-token chunks, ~280 ms steps), then reconsider the Flash-Next default
  • Re-test the eager islands inside F&P (GDN full-layer op, FA-RDNA2 op, QSA op) now that the replay wrapper is gone; inactive in all validated recipes
  • Make the opt-in HIP prefill paths faster than the defaults (all ported from 0.28, all slower today): QSA HIP (VLLM_RDNA_QSA_HIP, −4 to −10 % vs Triton), GDN HIP prefill chain (VLLM_GDN_HIP_PREFILL, −5.5 % vs FLA, delta_h/o unit tests failing; pairs with the delta_h LDS backport)
  • Flash-Next MoE W4A8 serve validation (kernel ported + unit-tested 27/27; stays opt-in)
  • Fix the 2 pre-existing unit failures test_mamba_align_chunk_split[3200] and test_gpu_block_table (failing since the port, found in the final unit run)
  • TP=8 across both PLX switches (never run; not validated in this PR)
  • W8A16 MoE: included, untested (needs a checkpoint); test/fix later
  • GDN delta_h LDS staging (v0.29/v0.30 92e8924078; 10-30 % kernel, only matters with the opt-in HIP GDN prefill chain)
  • FA-RDNA2 head-72 ViT prefill kernel (legacy fork 9406cd979; 1.8-3.0× vs Triton; would replace the SDPA fallback behind the vision max_pixels stopgap; needs varlen)
  • Wire the FA-RDNA2 int8/fp8 KV cache (humanwork-legacy 3baecdb516; A/B upstream int8_per_token_head first)
  • Cap FLA Triton autotune (num_warps ≤ 4) on gfx10 (legacy-upstream a1ae6bd7a; cold-start autotune hangs): investigate
  • RDNA3 attention tile tuning + int4/int8 KV prefill (jartx-rdna3 502e468)

Test plan (8× V620)

  • Editable build (venv-7.14.0, PYTORCH_ROCM_ARCH=gfx1030) and _rocm_C op registration check
  • pre-commit + mypy on fork-owned files (dc8083dbf5..ab5a35deed): all hooks pass, mypy clean on Linux; ~20 undocumented debug dumps removed, kept debug aids go through the logger. Serve-checked on EXL3 MTP0 and Flash-Next MTP2 (probes/prefix/after-load PASS; Flash-Next MTP2 1k c=8 194.3, 16k c=8 92.4 tok/s)
  • Unit tests: RDNA2 MoE/W4A8/MXFP4/EXL3 kernels, tests/models/qwen4_exp/, tests/v1/core/, GDN metadata, block table, prefill cadence, offloading connector, ple_offload (rdna_ar GPU test added). Final unit run on f419d8d6c0: 45 files pass. Known failures, all present before tonight's work (checked by running the same tests on the 00fcc97ff5 Python tree): opt-in HIP GDN prefill (delta_h 29 / o 31 cases), rocm_skinny_gemms abort, fused shared expert (12), part of compile/test_config (13), test_mamba_align_chunk_split (1 case: Flash-Next TP=4 replay-state at prompt 3200 ends a chunk at 3199), test_gpu_block_table (1 case: padded rows not -1); the last two are in the backlog
  • Flash-Next AWQ W4A16, TP=4 on one PLX, fp16:
    • Boots with upstream's PLE offload: memory and coherence
    • 8×16k/1k c=8 vs 74.2 tok/s: 81.6 (MTP0, with rdna_ar fix)
    • Correctness probes (before and after load, MTP0 and MTP2)
    • Prefix hits: repeat-prompt probe passes (MTP0 and MTP2)
    • MTP=2: ≈ 55 and 127 tok/s → 88.3 (16k/1k c=8) and 188.1 (1k/512 c=8)
    • F&P capture: no dmesg faults
    • RDNA2 MoE backend selected
  • Qwen3.8-27B AWQ dense with rdna_ar one-shot active (after the GDN norm fix v0.31 beats 0.28 on all 4 cells: 1k c=1 40.7 vs 39.7, 16k c=8 45.2 vs 41.3 tok/s)
  • Qwen3.5/3.6 GDN hybrid: Qwen3.6-35B-A3B GPTQ W4A16 on F&P (MTP0 + MTP=2) and PIECEWISE-only, all gates PASS on f419d8d6c0 (PIECEWISE concurrency flag = solo-vs-solo noise). W8A16 MoE: included, untested (no checkpoint), backlog
  • EXL3 27B, TP=4, MTP=2 (F&P; at parity with 0.28 when both benched greedy; MTP=2 c=8 step time ~470 ms on both trees is the next perf lever)
  • W4A8 on 27B AWQ (VLLM_RDNA2_W4A8_SDOT4=1, opt-in, prefill M ≥ 33): greedy output byte-identical to off, PPL 4.90 vs 4.94, prefix PASS, 0 faults; fast path fires on all 4 ranks for all 5 dense shapes. Prefill +24-34 % (table above), decode at parity. Unit tests 18/18 + 27/27 (b03bf9d210 fixes a stale assertion: the M ≤ 32 W4A16 prefill arm is not run-to-run deterministic). Still opt-in (default unchanged). Flash-Next/MoE W4A8 in the backlog
  • DeepSeek-V4-Flash MXFP4 (MXF4_RDNA2 MoE, RDNA2 sparse-MLA/indexer HIP kernels, fp16 attention): F&P TP=4, PPL 8.31-8.54, probe 3/3, prefix PASS, soaked (5 boots, 480 requests, 0 hangs); 4 cells in the results table (1k c=1 decode 38.8 tok/s/req, 16k prefill 1092 tok/s). Pruned checkpoint: PPL probe is the gate
  • Prefill cadence / adaptive prefill vs upstream adaptive LPT — superseded by the mixed-batch stall work: stall probe + gates, decode-stall cap (on for dense, opt-in for Flash-Next)
  • Concurrency correctness gate (concurrency_probe.py, on by default in serve_validate.sh): PASS on every serving config of the final test plan
  • TunableOp coverage: one profile folder rocm7.14-rocblas5.5 with 1601 rows/rank covering Flash-Next, 27B AWQ (W4A8 off/on), 27B EXL3 MTP0/MTP2, Qwen3.6 MTP0/MTP2, DSV4; lookup-hit 1601/1601
  • Final full test plan on f419d8d6c0: 15 serve runs, all gates PASS, 0 faults (table above)
  • Unit suite on the tip (ed4a1c5ea7 list, 2026-10-11): all fork files pass incl. the new test_hc_prefill_rdna2 34/34, test_decode_stall_bound 16/16, test_stall_probe 7/7; failures are the known pre-existing ones (opt-in HIP GDN prefill chain kkt/delta_h/o, test_rocm_skinny_gemms abort, test_fused_shared_expert 12, compile/test_config 12, test_mamba_align_chunk_split[3200], test_gpu_block_table) plus test_gdn_metadata_builder::test_decode_arena_max_bs_covers_capture_size, a stale expectation since ad356d1135 (arenas sized by the decode batch; the V2 runner never captures decode graphs above max_num_seqs): test updated (076fd4d636), re-run 53/53. test_rdna_ar_gpu skips with one visible GPU
  • Final pass on 3b43339ab0 → 20edef5321 (fused hyper-connection prefill, stall-cap fit per model, replay-copy wrapper removed, breakable strong refs, recipes): all configs PASS; full.env fixed in dc3065852b (cap off at TP=2), agentic on the mtp0 block, 27B stall cap 2672 → 320 ms
  • TP=8 across both PLX switches — moved to the backlog (not validated in this PR)

BlivionIaG and others added 30 commits October 7, 2026 23:46
@BlivionIaG @GeorgeMurAlkh @leapdragon
build(rocm): gfx1030 build plumbing on v0.31.0
a618ab3 
gfx1030 in the ROCm base image arch list and the RDNA2 kernel sources /
compile definitions in CMake (VLLM_ROCM_GFX1030, EXL3, rdna_allreduce).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 2 file(s):
  923672d 2026-07-12 BlivionIaG: rdna2 w4a16: add multi-config prefill kernel + 3-bucket dispatcher
  b1b3fa9 2026-07-17 BlivionIaG: Add w4a16 rdna2 moe + test + oom fix
  7c3fec8 2026-07-31 BlivionIaG: feat(rocm): enable gfx1030 (RDNA2) as a recognized RDNA arch
  c977282 2026-08-09 BlivionIaG: feat(rocm): paged MQA logits HIP kernel for DeepSeek V4 Lightning Indexer on gfx1030
  3173867 2026-08-13 BlivionIaG: feat(rocm): add native RDNA2 W8A8-FP8 dense GEMM for DeepSeek V4 Flash
  b55ccab 2026-08-13 BlivionIaG: feat(rocm): add native RDNA2 W4A4 MXFP4 fused MoE + dense GEMM kernels for gfx1030
  ea4e61a 2026-08-23 BlivionIaG: feat: hand-port GDN chunked prefill to HIP RDNA2 (5 kernels)
  5552701 2026-08-23 BlivionIaG: feat: hand-port GDN packed decode to HIP RDNA2
  26f329a 2026-08-27 BlivionIaG: wip(rocm): exl3 RDNA2 dot2 kernel - wire ops and refine kernels
  fc9012c 2026-08-31 BlivionIaG: rdna: exl3 6bpw kernel + one-time dequant for lm_head
  add108f 2026-09-02 BlivionIaG: rdna: define VLLM_ROCM_GFX1030 for _rocm_C target (schema registration bug)
  6920018 2026-09-02 BlivionIaG: rdna: fix two Phase4b-3 build errors
  5206533 2026-09-02 BlivionIaG: rdna: port K192 RDNA2 kernel work (W8A16 + W8A16-FP8 + sparse MLA)
  632dea3 2026-09-03 BlivionIaG: feat(rocm): register missing RDNA2 kernel ops in _rocm_C
  feb7b45 2026-09-06 BlivionIaG: feat: add exllama-clone AWQ prefill kernel for high-M (gfx1030)
  a406064 2026-09-08 BlivionIaG: [review-only] cherry-pick leapdragon rdna_ar fabric knobs onto rdna_extras (#1)
  5eb84b4 2026-09-08 BlivionIaG: feat(rocm): AWQ prefill kernel + causal_conv1d HIP + platform gates
  1046782 2026-09-11 BlivionIaG: cleanup: remove dead AWQ prefill kernel and binding
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  6635424 2026-09-13 BlivionIaG: build(rocm): add mrope_rdna2.cu to the gfx1030 EXT_SRC list
  1eb11e1 2026-09-14 BlivionIaG: feat(rocm): opt-in HIP paths for Qwen4Exp / Qwen3.8-Flash-Next on gfx1030
  0821be8 2026-09-22 George Muravei-Alkhavoi: Port opt-in served resident skinny MoE decode
  4d94aaf 2026-09-28 BlivionIaG: feat(rocm): wire W4A8 (int4 x int8 sdot4) as an opt-in dense W4A16 prefill replacement
  c9122de 2026-09-29 BlivionIaG: feat(rocm): fused MoE W4A8 sdot4 for gfx1030 (Flash-Next reachable)

Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
@GeorgeMurAlkh @BlivionIaG
Port opt-in served resident skinny MoE decode
0994bce 
gfx1030 in the ROCm base image arch list and the RDNA2 kernel sources /
compile definitions in CMake (VLLM_ROCM_GFX1030, EXL3, rdna_allreduce).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 1 file(s):
  0821be8 2026-09-22 George Muravei-Alkhavoi: Port opt-in served resident skinny MoE decode
@BlivionIaG @leapdragon
@GeorgeMurAlkh @eastwood-c
feat(rocm): RDNA2 HIP kernel library and op bindings
00ac7ee 
csrc/rocm RDNA2 kernels (W4A16/W4A8/W8A16/MXFP4 dense+MoE, FA-RDNA2,
GDN prefill/decode, causal_conv1d, QSA, HC, PLE conv, sparse MLA,
rdna_allreduce) and their torch op bindings / Python wrappers.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 39 file(s):
  fabf514 2026-06-11 BlivionIaG: Add Dense W4A16 RDNA2 kernel port
  923672d 2026-07-12 BlivionIaG: rdna2 w4a16: add multi-config prefill kernel + 3-bucket dispatcher
  b1b3fa9 2026-07-17 BlivionIaG: Add w4a16 rdna2 moe + test + oom fix
  c977282 2026-08-09 BlivionIaG: feat(rocm): paged MQA logits HIP kernel for DeepSeek V4 Lightning Indexer on gfx1030
  3173867 2026-08-13 BlivionIaG: feat(rocm): add native RDNA2 W8A8-FP8 dense GEMM for DeepSeek V4 Flash
  b55ccab 2026-08-13 BlivionIaG: feat(rocm): add native RDNA2 W4A4 MXFP4 fused MoE + dense GEMM kernels for gfx1030
  ea4e61a 2026-08-23 BlivionIaG: feat: hand-port GDN chunked prefill to HIP RDNA2 (5 kernels)
  5552701 2026-08-23 BlivionIaG: feat: hand-port GDN packed decode to HIP RDNA2
  26f329a 2026-08-27 BlivionIaG: wip(rocm): exl3 RDNA2 dot2 kernel - wire ops and refine kernels
  567f243 2026-08-31 BlivionIaG: rdna: drop redundant size_m/n/k from exl3_gemm_rdna2 — fix dynamo
  7e18f20 2026-08-31 BlivionIaG: rdna: fix EXL3 allow_in_graph decorator — use torch._dynamo not torch.compiler
  20679f5 2026-08-31 BlivionIaG: rdna: allow_in_graph EXL3 op wrappers to fix dynamo shape guard
  fc9012c 2026-08-31 BlivionIaG: rdna: exl3 6bpw kernel + one-time dequant for lm_head
  719c7b0 2026-09-02 BlivionIaG: feat(exl3): add Python wrapper for exl3_dequant_bits6_mul1
  ec645bd 2026-09-02 BlivionIaG: rdna: backport 14 missing ops.h RDNA2 declarations
  4fa410f 2026-09-02 BlivionIaG: rdna: fix e2m1_lut consumer calls (use function syntax, not subscript)
  9b83b95 2026-09-02 BlivionIaG: rdna: convert e2m1_lut to a __device__ switch function (works on gfx1030)
  b513bdd 2026-09-02 BlivionIaG: rdna: fix e2m1_lut dynamic-init (plain constexpr, not __device__ constexpr)
  6920018 2026-09-02 BlivionIaG: rdna: fix two Phase4b-3 build errors
  5206533 2026-09-02 BlivionIaG: rdna: port K192 RDNA2 kernel work (W8A16 + W8A16-FP8 + sparse MLA)
  6a285cf 2026-09-02 BlivionIaG: rdna: fix int32 overflow in paged_mqa slot_base on strided indexer layout
  faf87f8 2026-09-03 BlivionIaG: fix(exl3): mark exl3_hadamard_128 output as mutating in op schema
  632dea3 2026-09-03 BlivionIaG: feat(rocm): register missing RDNA2 kernel ops in _rocm_C
  e85909f 2026-09-03 BlivionIaG: feat(attention): add standalone RDNA_ATTN backend with proven FA-RDNA2 kernels
  63584d6 2026-09-04 BlivionIaG: perf(attn): GQA-aware D=256 paged decode kernel (1.5-2.0x FA decode)
  aac1fcd 2026-09-04 BlivionIaG: fix(gdn): chunk-local indexing bug in delta_h prefill kernel
  62694f2 2026-09-04 BlivionIaG: feat(exl3): decode-trellis prefill path (7.5x faster prefill GEMM)
  feb7b45 2026-09-06 BlivionIaG: feat: add exllama-clone AWQ prefill kernel for high-M (gfx1030)
  cafe95e 2026-09-08 BlivionIaG: fix(rocm): FPP13 checkpoint — FULL_AND_PIECEWISE greedy is coherent
  a406064 2026-09-08 BlivionIaG: [review-only] cherry-pick leapdragon rdna_ar fabric knobs onto rdna_extras (#1)
  7779514 2026-09-08 BlivionIaG: fix(rocm,fa-rdna2): torch::zeros for O_partial/M_partial on RDNA2 cudagraph replay
  bd9fc90 2026-09-08 BlivionIaG: fix(rocm,fa-rdna2): torch::zeros for output O on RDNA2 cudagraph replay
  5eb84b4 2026-09-08 BlivionIaG: feat(rocm): AWQ prefill kernel + causal_conv1d HIP + platform gates
  3326f03 2026-09-08 BlivionIaG: fix(rocm,gdn-decode): honor NULL_BLOCK_ID=0 sentinel in decode kernel
  46fccc1 2026-09-08 BlivionIaG: test(gdn): extend GDN prefill regression test to multi-chunk + non-aligned T
  c091c42 2026-09-10 BlivionIaG: wip: checkpoint-validated state + chase csrc + debug instruments + row-zero fix
  13a3daf 2026-09-11 BlivionIaG: cleanup(gqa): drop true-GQA variant, keep subgroup + off
  ecfec4e 2026-09-11 BlivionIaG: perf: GQA multi-head FA-RDNA2 prefill kernels for gfx1030
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP
  71a5455 2026-09-11 BlivionIaG: perf: HIP gated_rms_norm kernel for Qwen3.x GDN RMSNormGated
  56f6711 2026-09-11 BlivionIaG: cleanup(w4a16): drop measured-slower experiment configs, keep debug knobs
  441ec05 2026-09-11 BlivionIaG: refactor(w4a16): generalize prefill inner loop, add dormant ConfigA_Large slot
  7ac98a2 2026-09-11 BlivionIaG: fix: revert ConfigH (K_STEP=64) — broken for M > 256, keep ConfigA
  1046782 2026-09-11 BlivionIaG: cleanup: remove dead AWQ prefill kernel and binding
  aaaae85 2026-09-11 BlivionIaG: perf: route W4A16 AWQ prefill through GPTQ ConfigA kernel
  854dd56 2026-09-12 BlivionIaG: fix(rocm): sanitize GDN state caches under cudagraph capture
  c6b5cfb 2026-09-12 BlivionIaG: fix(rocm): make gptq_gemm_rdna2 activation contiguous-safe
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  bb48489 2026-09-13 BlivionIaG: fix(rocm): move mrope_forward_rdna2 host wrapper to global scope
  4f1ee27 2026-09-13 BlivionIaG: fix(rocm): add missing c10/cuda/CUDAGuard.h include to mrope_rdna2.cu
  cf055ad 2026-09-13 BlivionIaG: feat(rocm): HIP M-RoPE forward (mrope_forward_rdna2)
  82b6f18 2026-09-13 BlivionIaG: fix(rocm): enable HIP causal_conv1d_fwd_rdna2 (out-stride + null-block)
  1eb11e1 2026-09-14 BlivionIaG: feat(rocm): opt-in HIP paths for Qwen4Exp / Qwen3.8-Flash-Next on gfx1030
  b26763e 2026-09-15 BlivionIaG: fix(rocm): ple_short_conv optional-tensor schema + on_gfx10x import path
  02adbfd 2026-09-15 BlivionIaG: feat(rocm): fp16 state support in gdn_decode_rdna2
  8960a3b 2026-09-16 BlivionIaG: fix(rocm): make the Qwen4Exp HC HIP chain compute correctly
  d0d577f 2026-09-16 BlivionIaG: fix(rocm): add the 10 missing Qwen4Exp _custom_ops wrappers
  c350fa2 2026-09-16 BlivionIaG: merge(rdna2): port PR #11 wvSplitK ownership/decode + rdna_ar device fix
  d1b200b 2026-09-17 BlivionIaG: perf(rocm): cut FA-RDNA2 GQA prefill cost ~2x with a register-resident O accumulator
  3b59ee1 2026-09-18 BlivionIaG: feat(rocm): opt-in leap T44b rdna_ar (VRAM flags + wedge check) (#13)
  0821be8 2026-09-22 George Muravei-Alkhavoi: Port opt-in served resident skinny MoE decode
  6ca0fc6 2026-09-27 BlivionIaG: bench(rocm): RDNA_ATTN split-decode A/B for verify and mixed batches
  9dce178 2026-09-27 BlivionIaG: perf(rocm): FA-RDNA2 decode kernel for spec-decode verify and mixed batches
  690cbf0 2026-09-27 BlivionIaG: perf(rocm): FA-RDNA2 causal tile skipping, in-place output, parallel GQA softmax
  cf20bd2 2026-09-28 BlivionIaG: fix(rocm): W4A8 ops return Tensors (empty = not eligible), not int64 status
  b95036b 2026-09-28 BlivionIaG: feat(rocm): 27B dense launcher + keep the W4A8 attempt out of dynamo's graph
  4d94aaf 2026-09-28 BlivionIaG: feat(rocm): wire W4A8 (int4 x int8 sdot4) as an opt-in dense W4A16 prefill replacement
  c9122de 2026-09-29 BlivionIaG: feat(rocm): fused MoE W4A8 sdot4 for gfx1030 (Flash-Next reachable)
  3050090 2026-09-29 BlivionIaG: fix(rdna2): preserve pre-W4A8 W4A16 split-K choice; W4A8 defaults to 0
  a641d65 2026-09-29 BlivionIaG: fix(rdna2): K_STEP-align prefill split-K (W4A8 fallback NaN) + W4A8 selector guard
  45e218e 2026-09-29 BlivionIaG: feat(rdna2): env-gated W4A8 shape diagnostic + 27B dense A/B tooling
  2138382 2026-09-29 BlivionIaG: feat(rocm): proper vLLM-idiomatic W4A8 prefill integration
  8646106 2026-09-30 BlivionIaG: feat(rdna2): canonical VLLM_RDNA2_MOE_FP32_ACCUM env + per-call fp32_accum
  09b7f33 2026-09-30 BlivionIaG: feat(rdna2): fp32 atomic accum default for MoE epilogue (gfx1030)
  adcd0f2 2026-09-30 BlivionIaG: bench(explore): PR30 harness fixes + V620 results (H1 holds, H3 fails)
  4b37b52 2026-09-30 BlivionIaG: docs(explore): RDNA2 W4A16 M-dispatch under torch.compile
  7be99d8 2026-10-01 BlivionIaG: feat(rdna): add exl3_project_rdna2 fused decode projection op
  9cad879 2026-10-01 BlivionIaG: [rdna2][MoE-epilogue] default back to packed-fp16 CAS; fp32 scratch opt-in
  c5fd67e 2026-10-02 BlivionIaG: fix(rdna2): respect row/cache width strides in qsa_store_cache_rows_kernel
  d19b87e 2026-10-02 BlivionIaG: port: QSA RDNA2 HIP decode path (v0.30 fix files onto rdna_extras)
  284fbad 2026-10-02 BlivionIaG: [FA-RDNA2][rdna2] gqa prefill kernel templated on HEAD_DIM (128/256)
  c6d99ca 2026-10-06 BlivionIaG: fix(rdna2): HC HIP stream correctness + MoE prefill tile

Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
@leapdragon @BlivionIaG @eastwood-c
feat(rocm): port Aron Hsiao's kernels-csrc changes (3 commits)
cd100da 
csrc/rocm RDNA2 kernels (W4A16/W4A8/W8A16/MXFP4 dense+MoE, FA-RDNA2,
GDN prefill/decode, causal_conv1d, QSA, HC, PLE conv, sparse MLA,
rdna_allreduce) and their torch op bindings / Python wrappers.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 2 file(s):
  c091c42 2026-09-10 BlivionIaG: wip: checkpoint-validated state + chase csrc + debug instruments + row-zero fix
  c05af40 2026-09-11 Aron Hsiao: feat(rdna2): ungate Hybrid W4A16 on gfx10, keep RDNA2 auto default
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
@BlivionIaG @GeorgeMurAlkh
@leapdragon @wsantos @eastwood-c
feat(rocm): RDNA2 platform detection, config and env knobs
9f3dc48 
RDNA platform detection, gfx10x helpers, env knobs and config
plumbing (adaptive prefill, breakable-cudagraph archs, ...).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 13 file(s):
  fabf514 2026-06-11 BlivionIaG: Add Dense W4A16 RDNA2 kernel port
  3c85702 2026-07-15 BlivionIaG: feat(rocm): Wire FA-RDNA2 dispatch into RocmAttentionImpl
  7c3fec8 2026-07-31 BlivionIaG: feat(rocm): enable gfx1030 (RDNA2) as a recognized RDNA arch
  70e65f9 2026-08-27 BlivionIaG: fix(distributed): execute TP collectives under torch.compile on ROCm
  0193296 2026-08-28 BlivionIaG: wip(rocm): exl3 mul1 codebook markers + folded-weight cache
  f8722ea 2026-08-28 BlivionIaG: wip(rocm): exl3 quant method registration (Exl3Config) for RDNA
  dda2a1d 2026-09-02 BlivionIaG: rdna: backport RDNA detection gates from rdna2_extras to vllm/platforms/rocm.py
  ea78104 2026-09-03 BlivionIaG: fix(rocm): platform gates + MoE W8A16 parity (completes fb1687b)
  fb1687b 2026-09-03 BlivionIaG: feat(rocm): VLLM_FORCE_CUSTOM_ALL_REDUCE for PCIe-only RDNA
  e85909f 2026-09-03 BlivionIaG: feat(attention): add standalone RDNA_ATTN backend with proven FA-RDNA2 kernels
  63c2ed5 2026-09-05 BlivionIaG: feat(env): register DBG_VLLM_STEP_TIMING for worker propagation
  cafe95e 2026-09-08 BlivionIaG: fix(rocm): FPP13 checkpoint — FULL_AND_PIECEWISE greedy is coherent
  5eb84b4 2026-09-08 BlivionIaG: feat(rocm): AWQ prefill kernel + causal_conv1d HIP + platform gates
  631aa35 2026-09-08 BlivionIaG: feat(rocm,gdn): eager_break_during_capture for Qwen3.5/3.8 GDN + architecture whitelist
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  965c3b9 2026-09-13 BlivionIaG: feat(rdna): add PIX topology helpers without changing AR policy
  1c1dbee 2026-09-13 BlivionIaG: fix(rocm): disable VLLM_ROCM_MOE_SKINNY by default (recipe 0008 kernel is wrong)
  cf055ad 2026-09-13 BlivionIaG: feat(rocm): HIP M-RoPE forward (mrope_forward_rdna2)
  1eb11e1 2026-09-14 BlivionIaG: feat(rocm): opt-in HIP paths for Qwen4Exp / Qwen3.8-Flash-Next on gfx1030
  59237b3 2026-09-14 BlivionIaG: fix(rdna2): repair 0008 skinny default, gfx10 bf16 abort, warmup timeout
  9c9509b 2026-09-16 BlivionIaG: Revert "feat(rocm): default the Qwen4Exp HIP kernels on (tracker S6)"
  cb0d441 2026-09-16 BlivionIaG: feat(rocm): default the Qwen4Exp HIP kernels on (tracker S6)
  c350fa2 2026-09-16 BlivionIaG: merge(rdna2): port PR #11 wvSplitK ownership/decode + rdna_ar device fix
  c59a23f 2026-09-17 BlivionIaG: chore: recover remote-only work products into the repo
  5c4ab98 2026-09-17 BlivionIaG: revert(rocm): drop the PR #5 wvSplitK gfx1030 decode port
  b33f9b6 2026-09-18 BlivionIaG: fix(rocm): fall back to torch for get_device_name when amdsmi sees no GPUs
  4425834 2026-09-18 BlivionIaG: fix(rocm): make platform detection and worker init robust on RDNA consumer GPUs
  16695b2 2026-09-18 BlivionIaG: test(rocm): isolate V620 amdsmi and PLE fp8 CPU tests from HIP
  b8354ff 2026-09-18 BlivionIaG: feat(rocm): V620 MoE JSON, ROCR amdsmi index, PLE fp8 fused gather
  3b59ee1 2026-09-18 BlivionIaG: feat(rocm): opt-in leap T44b rdna_ar (VRAM flags + wedge check) (#13)
  0821be8 2026-09-22 George Muravei-Alkhavoi: Port opt-in served resident skinny MoE decode
  9a634c7 2026-09-22 George Muravei-Alkhavoi: Port resident W4A16 MoE and FP16 shared experts to current base
  cceabbd 2026-09-23 Waldecir Santos: fix(core): keep FULL decode graphs when piecewise capture is unavailable
  295a842 2026-09-24 BlivionIaG: feat(rocm): make rdna_ar cover the custom-allreduce cases
  48c56ef 2026-09-24 BlivionIaG: fix(rdna2): make VLLM_USE_RDNA2_FA=1 actually select RDNA_ATTN
  6c26c78 2026-09-24 BlivionIaG: fix(rocm): keep FULL decode graphs alongside piecewise capture
  700753d 2026-09-25 BlivionIaG: fix(rocm,mtp): unbreak MTP draft inference — amdsmi shutdown masking + local-argmax draft sampling
  97fbae5 2026-09-27 BlivionIaG: [ROCm][RDNA] Load and call hippihx V1 from RDNA_ATTN decode (opt-in)
  22d6e34 2026-09-27 BlivionIaG: feat(rdna2): two-shot rdna_ar (PR #22) + wide gate default, validated on 4x V620
  f22fee0 2026-09-27 George Muravei-Alkhavoi: perf(rocm): opt in to HIP MoE for concurrent MTP decode
  6ed3909 2026-09-28 BlivionIaG: feat(engine): prefill step cadence for non-DP engines
  8646106 2026-09-30 BlivionIaG: feat(rdna2): canonical VLLM_RDNA2_MOE_FP32_ACCUM env + per-call fp32_accum
  9cad879 2026-10-01 BlivionIaG: [rdna2][MoE-epilogue] default back to packed-fp16 CAS; fp32 scratch opt-in

Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
Co-authored-by: Waldecir Santos <waldecir@gmail.com>
Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
@leapdragon @BlivionIaG
feat(rocm): port Aron Hsiao's platform-config changes (2 commits)
ea4b784 
RDNA platform detection, gfx10x helpers, env knobs and config
plumbing (adaptive prefill, breakable-cudagraph archs, ...).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 4 file(s):
  c05af40 2026-09-11 Aron Hsiao: feat(rdna2): ungate Hybrid W4A16 on gfx10, keep RDNA2 auto default
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
@BlivionIaG @leapdragon
feat(rdna2): W4A16/W4A8/W8A16/FP8 dense linear paths
ccc8ad0 
RDNA2 dense linear kernels in the mixed-precision registry and
layers (W4A16 RDNA2/hybrid, W4A8, W8A16, int8 dense shadow).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 22 file(s):
  fabf514 2026-06-11 BlivionIaG: Add Dense W4A16 RDNA2 kernel port
  923672d 2026-07-12 BlivionIaG: rdna2 w4a16: add multi-config prefill kernel + 3-bucket dispatcher
  7c3fec8 2026-07-31 BlivionIaG: feat(rocm): enable gfx1030 (RDNA2) as a recognized RDNA arch
  b55ccab 2026-08-13 BlivionIaG: feat(rocm): add native RDNA2 W4A4 MXFP4 fused MoE + dense GEMM kernels for gfx1030
  ea4e61a 2026-08-23 BlivionIaG: feat: hand-port GDN chunked prefill to HIP RDNA2 (5 kernels)
  5552701 2026-08-23 BlivionIaG: feat: hand-port GDN packed decode to HIP RDNA2
  f8722ea 2026-08-28 BlivionIaG: wip(rocm): exl3 quant method registration (Exl3Config) for RDNA
  45b5be4 2026-09-02 BlivionIaG: rdna: restore upstream test_w4a16_kernel_selection.py (PR vllm-project#40977)
  53e2f2c 2026-09-02 BlivionIaG: rdna: fix gfx1030_only skipif (hasattr returns False even when op is registered)
  5206533 2026-09-02 BlivionIaG: rdna: port K192 RDNA2 kernel work (W8A16 + W8A16-FP8 + sparse MLA)
  feb7b45 2026-09-06 BlivionIaG: feat: add exllama-clone AWQ prefill kernel for high-M (gfx1030)
  0ac2b30 2026-09-06 BlivionIaG: feat(rocm): add HIP AOT RMSNorm forward_rocm to plain RMSNorm
  d0bc0a2 2026-09-06 BlivionIaG: feat(rocm): port dense AWQ support to RDNA2 W4A16 kernel
  cafe95e 2026-09-08 BlivionIaG: fix(rocm): FPP13 checkpoint — FULL_AND_PIECEWISE greedy is coherent
  c05af40 2026-09-11 Aron Hsiao: feat(rdna2): ungate Hybrid W4A16 on gfx10, keep RDNA2 auto default
  71a5455 2026-09-11 BlivionIaG: perf: HIP gated_rms_norm kernel for Qwen3.x GDN RMSNormGated
  55ff004 2026-09-11 BlivionIaG: perf: fold Gemma (1+w) scale in x.dtype so vllm_c rms_norm fires
  aaaae85 2026-09-11 BlivionIaG: perf: route W4A16 AWQ prefill through GPTQ ConfigA kernel
  b205fc0 2026-09-12 BlivionIaG: debug(rocm): env-gated forensics probes for TP=4 decode NaN
  de8a4e4 2026-09-12 BlivionIaG: fix(rocm): disable RDNA2 W4A16 kernel at TP>2 (cudagraph decode corruption)
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  f86faad 2026-09-13 BlivionIaG: fix(rocm): enable RDNA2 W4A16 at TP=4 via breakable cudagraphs + PYNCCL
  c350fa2 2026-09-16 BlivionIaG: merge(rdna2): port PR #11 wvSplitK ownership/decode + rdna_ar device fix
  c59a23f 2026-09-17 BlivionIaG: chore: recover remote-only work products into the repo
  5c4ab98 2026-09-17 BlivionIaG: revert(rocm): drop the PR #5 wvSplitK gfx1030 decode port
  f0b4273 2026-09-28 BlivionIaG: fix(rocm): treat the W4A8 sentinel as None too (undefined at::Tensor -> Python None)
  933e3b1 2026-09-28 BlivionIaG: test(rocm): use the aq_ret name at the bad-group site
  75bc07d 2026-09-28 BlivionIaG: test(rocm): bad-group-size case expects the empty-tensor sentinel (8-space site)
  cf20bd2 2026-09-28 BlivionIaG: fix(rocm): W4A8 ops return Tensors (empty = not eligible), not int64 status
  158a9ad 2026-09-28 BlivionIaG: revert(rocm): drop @torch._dynamo.disable on the W4A8 helper
  b95036b 2026-09-28 BlivionIaG: feat(rocm): 27B dense launcher + keep the W4A8 attempt out of dynamo's graph
  e812092 2026-09-28 BlivionIaG: fix(rocm): drop the function-local 'import os' in the decode branch
  80ba0e3 2026-09-28 BlivionIaG: fix(rocm): restore the pre-existing W4A16 dispatch chain verbatim
  ace9fbf 2026-09-28 BlivionIaG: feat(rocm): log once when the W4A8 prefill path is taken
  4d94aaf 2026-09-28 BlivionIaG: feat(rocm): wire W4A8 (int4 x int8 sdot4) as an opt-in dense W4A16 prefill replacement
  a641d65 2026-09-29 BlivionIaG: fix(rdna2): K_STEP-align prefill split-K (W4A8 fallback NaN) + W4A8 selector guard
  2138382 2026-09-29 BlivionIaG: feat(rocm): proper vLLM-idiomatic W4A8 prefill integration
  956dbd4 2026-09-29 BlivionIaG: fix(rocm): W4A8 helper truly branch-free + shape pre-check (the earlier patches had silently missed)
  b1a06b3 2026-09-29 BlivionIaG: fix(rocm): make the W4A8 helper branch-free (dynamo inlines it into the traced forward)
  be22e16 2026-09-29 BlivionIaG: fix(rocm): W4A8 eligibility decided before the call (no data-dependent branch)
  adcd0f2 2026-09-30 BlivionIaG: bench(explore): PR30 harness fixes + V620 results (H1 holds, H3 fails)
  4b37b52 2026-09-30 BlivionIaG: docs(explore): RDNA2 W4A16 M-dispatch under torch.compile
  74315f4 2026-10-01 BlivionIaG: fix(rocm): gemma_rms_norm fake kernel contiguous output + exl3 27B mul1 TP=4 validation

Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
@leapdragon @BlivionIaG @GeorgeMurAlkh
feat(rdna2): port Aron Hsiao's quant-linear changes (2 commits)
a567946 
RDNA2 dense linear kernels in the mixed-precision registry and
layers (W4A16 RDNA2/hybrid, W4A8, W8A16, int8 dense shadow).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 7 file(s):
  fabf514 2026-06-11 BlivionIaG: Add Dense W4A16 RDNA2 kernel port
  5eea5c9 2026-08-28 BlivionIaG: wip(rocm): exl3 3inst decode working end-to-end (Qwen3.5-0.8B)
  c05af40 2026-09-11 Aron Hsiao: feat(rdna2): ungate Hybrid W4A16 on gfx10, keep RDNA2 auto default
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  59237b3 2026-09-14 BlivionIaG: fix(rdna2): repair 0008 skinny default, gfx10 bf16 abort, warmup timeout

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
@BlivionIaG @GeorgeMurAlkh @leapdragon
feat(rdna2): native RDNA2 MoE (W4A16/W4A8/W8A16/MXFP4) and dispatch
52e79c3 
RDNA2 fused-MoE experts and their dispatch: WNA16 oracle RDNA2_W4A16 backend,
compressed-tensors rocm_moe_rdna (RDNA2-only now that RDNA3 moved to the
upstream oracle, vllm-project#55522), MXF4_RDNA2, W8A16, resident decode, tuned configs.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 26 file(s):
  b1b3fa9 2026-07-17 BlivionIaG: Add w4a16 rdna2 moe + test + oom fix
  b55ccab 2026-08-13 BlivionIaG: feat(rocm): add native RDNA2 W4A4 MXFP4 fused MoE + dense GEMM kernels for gfx1030
  5206533 2026-09-02 BlivionIaG: rdna: port K192 RDNA2 kernel work (W8A16 + W8A16-FP8 + sparse MLA)
  ea78104 2026-09-03 BlivionIaG: fix(rocm): platform gates + MoE W8A16 parity (completes fb1687b)
  12c7702 2026-09-03 BlivionIaG: feat(moe): MXF4_RDNA2 backend routing + RDNA2Mxfp4MoEExperts
  8f578b1 2026-09-14 BlivionIaG: fix(rocm): match MoE quant-method subclasses in _needs_intermediate_size_param
  4b799bf 2026-09-15 BlivionIaG: fix(rocm): fix moe_align_block_size cudagraph race (8/8 concurrent coherence)
  c6cde0f 2026-09-15 BlivionIaG: fix(rocm): cache scales/zeros and pre-allocate buffers for RDNA2 W4A16 MoE
  b549c22 2026-09-15 BlivionIaG: feat(rocm): wire RDNA2 W4A16 MoE kernel into oracle for Flash-Next
  b8354ff 2026-09-18 BlivionIaG: feat(rocm): V620 MoE JSON, ROCR amdsmi index, PLE fp8 fused gather
  0821be8 2026-09-22 George Muravei-Alkhavoi: Port opt-in served resident skinny MoE decode
  178f107 2026-09-22 George Muravei-Alkhavoi: Compare qualified V620 MoE tile against existing tile
  85f22d5 2026-09-22 George Muravei-Alkhavoi: Qualify eight-row resident MoE prefill tile for V620
  9a634c7 2026-09-22 George Muravei-Alkhavoi: Port resident W4A16 MoE and FP16 shared experts to current base
  f125afc 2026-09-22 George Muravei-Alkhavoi: Fix RDNA2 MoE FP16 zero dequantization on current base
  c9122de 2026-09-29 BlivionIaG: feat(rocm): fused MoE W4A8 sdot4 for gfx1030 (Flash-Next reachable)
  e3ced14 2026-09-30 BlivionIaG: test(rdna2): fp32 accum stability, rel-L2 vs CAS, cudagraph capture
  09b7f33 2026-09-30 BlivionIaG: feat(rdna2): fp32 atomic accum default for MoE epilogue (gfx1030)
  9cad879 2026-10-01 BlivionIaG: [rdna2][MoE-epilogue] default back to packed-fp16 CAS; fp32 scratch opt-in
  c6d99ca 2026-10-06 BlivionIaG: fix(rdna2): HC HIP stream correctness + MoE prefill tile

Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
@GeorgeMurAlkh @BlivionIaG
feat(rdna2): port George Muravei-Alkhavoi's moe changes (4 commits)
96d3027 
RDNA2 fused-MoE experts and their dispatch: WNA16 oracle RDNA2_W4A16 backend,
compressed-tensors rocm_moe_rdna (RDNA2-only now that RDNA3 moved to the
upstream oracle, vllm-project#55522), MXF4_RDNA2, W8A16, resident decode, tuned configs.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 7 file(s):
  f123860 2026-09-03 BlivionIaG: fix(exl3): compute MLP activation in fp32 to avoid fp16 overflow
  0821be8 2026-09-22 George Muravei-Alkhavoi: Port opt-in served resident skinny MoE decode
  1e68a8b 2026-09-22 George Muravei-Alkhavoi: Version served V620 MoE tuning and launch settings
  9a634c7 2026-09-22 George Muravei-Alkhavoi: Port resident W4A16 MoE and FP16 shared experts to current base
  a25e881 2026-09-27 George Muravei-Alkhavoi: test(rocm): qualify concurrent MTP skinny MoE decode shapes
  c9122de 2026-09-29 BlivionIaG: feat(rocm): fused MoE W4A8 sdot4 for gfx1030 (Flash-Next reachable)

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
@leapdragon @GeorgeMurAlkh
@BlivionIaG @eastwood-c
feat(rdna2): port Aron Hsiao's moe changes (2 commits)
6f43b9e 
RDNA2 fused-MoE experts and their dispatch: WNA16 oracle RDNA2_W4A16 backend,
compressed-tensors rocm_moe_rdna (RDNA2-only now that RDNA3 moved to the
upstream oracle, vllm-project#55522), MXF4_RDNA2, W8A16, resident decode, tuned configs.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 4 file(s):
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  59237b3 2026-09-14 BlivionIaG: fix(rdna2): repair 0008 skinny default, gfx10 bf16 abort, warmup timeout
  b2584f0 2026-09-15 George Muravei-Alkhavoi: fix(rocm): restore V620 CPU PLE and MTP startup compatibility
  ed94e3f 2026-09-18 BlivionIaG: fix(rocm,moe,spec): Qwen4Exp MTP path prerequisites from the MTP=2 investigation
  f22fee0 2026-09-27 George Muravei-Alkhavoi: perf(rocm): opt in to HIP MoE for concurrent MTP decode
  6574a75 2026-09-27 George Muravei-Alkhavoi: test(rocm): expose missing concurrent MTP skinny dispatch
  a25e881 2026-09-27 George Muravei-Alkhavoi: test(rocm): qualify concurrent MTP skinny MoE decode shapes

Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: BlivionIaG <kev29lt@gmail.com>
Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
@BlivionIaG
feat(rdna): EXL3 quantization on RDNA
93df2d1 
EXL3 (QTIP) quantization on RDNA: config/linear method, kernels glue,
tests and quantizer tooling.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 68 file(s):
  26f329a 2026-08-27 BlivionIaG: wip(rocm): exl3 RDNA2 dot2 kernel - wire ops and refine kernels
  40850e6 2026-08-27 BlivionIaG: wip(rocm): exl3 RDNA2 dot2 kernel port (in-progress)
  5332d65 2026-08-28 BlivionIaG: fix(rocm): exl3 loaders handle MergedLinear partition + TP fallback
  bbdb15c 2026-08-28 BlivionIaG: fix(rocm): exl3 head_bits=0 fallback + float-bits early-return + TP loader
  c771a31 2026-08-28 BlivionIaG: fix(rocm): remove unused scaffold vars in exl3 dense kernel
  0193296 2026-08-28 BlivionIaG: wip(rocm): exl3 mul1 codebook markers + folded-weight cache
  5eea5c9 2026-08-28 BlivionIaG: wip(rocm): exl3 3inst decode working end-to-end (Qwen3.5-0.8B)
  f8722ea 2026-08-28 BlivionIaG: wip(rocm): exl3 quant method registration (Exl3Config) for RDNA
  640da7c 2026-08-29 BlivionIaG: fix(rocm): exl3 loader copy direction + free original fp16 weight
  bbcab03 2026-08-29 BlivionIaG: fix(rocm): contiguous output slice for fused-loop Hadamard
  b746b25 2026-08-29 BlivionIaG: feat(rocm): exl3 on-the-fly dequant for fused layers via kernel loop
  8f2583d 2026-08-31 BlivionIaG: Revert "rdna: disable dynamo on Exl3LinearMethod.apply"
  567f243 2026-08-31 BlivionIaG: rdna: drop redundant size_m/n/k from exl3_gemm_rdna2 — fix dynamo
  646d8ad 2026-08-31 BlivionIaG: rdna: disable dynamo on Exl3LinearMethod.apply
  609bd73 2026-08-31 BlivionIaG: rdna: exl3 dequant host wrapper to global scope
  fc9012c 2026-08-31 BlivionIaG: rdna: exl3 6bpw kernel + one-time dequant for lm_head
  126ecfa 2026-08-31 BlivionIaG: rdna: exl3 6bpw lm_head NotImplementedError + RDNA warmup guard
  a34a769 2026-08-31 BlivionIaG: rdna: exl3 per-partition runtime GEMM for MergedLinear (no exllamav3 needed)
  1ba85d2 2026-09-01 BlivionIaG: fix(exl3): free unused _w_fp16 buffers + use x.size(0) for dynamo
  29b7790 2026-09-01 BlivionIaG: fix(exl3): use layer._exl3_M_MAX in CG-PATH output padding (NameError)
  8465244 2026-09-01 BlivionIaG: fix(exl3): use full-size svh in FB-PATH when M > M_MAX
  8e59412 2026-09-01 BlivionIaG: fix(exl3): proper cudagraph fix — torch.zeros in create_weights + clean apply()
  b0d32a7 2026-09-01 BlivionIaG: fix(exl3): use .clone() instead of *1.0 for guaranteed page commit
  0a5a47f 2026-09-01 BlivionIaG: fix(exl3): commit GPU pages with x*1.0 to avoid uncommitted page fault
  69f4909 2026-09-01 BlivionIaG: fix(exl3): pad out_i in FB-PATH to match out shape for broadcast assignment
  0408d94 2026-09-01 BlivionIaG: fix(exl3): pad output back to _original_M for downstream shape assertions
  e24e0da 2026-09-01 BlivionIaG: fix(exl3): hard-cap M to max_num_batched_tokens to avoid uncommitted page fault
  dad9f8b 2026-09-01 BlivionIaG: fix(exl3): slice x to actual storage capacity — handle vLLM pre-allocated buffers
  c9d1d80 2026-09-01 BlivionIaG: debug(exl3): log x.shape/x.device/x.data_ptr in apply() when VLLM_EXL3_APPLY_DBG=1
  57dfe2d 2026-09-01 BlivionIaG: debug(exl3/hadamard): log input dims when rows>1024 or cols>16384
  f748edb 2026-09-01 BlivionIaG: fix(exl3): ensure svh is on x.device in FB-PATH fallback — fix Hadamard kernel host-page-fault
  ad17865 2026-09-01 BlivionIaG: feat(exl3): wrap kernels in @torch._dynamo.allow_in_graph
  07810e3 2026-09-01 BlivionIaG: feat(exl3): pre-slice svh per partition in process_weights_after_loading
  8523550 2026-09-01 BlivionIaG: fix(exl3): move trellis copy to process_weights_after_loading
  f722e6f 2026-09-01 BlivionIaG: feat(exl3): cudagraph-friendly buffer scaffolding (CG-PATH disabled)
  d184be5 2026-09-02 BlivionIaG: perf(exl3): gate Hadamard debug print on VLLM_EXL3_HADAMARD_DBG env var
  e78f528 2026-09-02 BlivionIaG: debug(exl3): log lm_head dequant norm/max to diagnose zero logits
  04208a5 2026-09-02 BlivionIaG: fix(exl3): copy dequantized weight into lm_head.weight Parameter in place
  ff6896c 2026-09-02 BlivionIaG: fix(exl3): swap lm_head.weight AFTER bits==6 dequant fills _w_fp16
  39897be 2026-09-02 BlivionIaG: fix(exl3): keep layer.weight alive when folded into _w_fp16 (lm_head)
  97346a7 2026-09-02 BlivionIaG: fix(exl3): derive k_tiles from layer.trellis.shape[0], not input_size/in_features
  e2a71c3 2026-09-02 BlivionIaG: fix(exl3): reallocate _exl3_bufs_trellis when bits != 3
  6955ce0 2026-09-02 BlivionIaG: fix(exl3): guard _exl3_log calls with is_compiling() to unblock dynamo trace
  9f98b43 2026-09-03 BlivionIaG: fix(exl3): fold mul1-marked layers to dense fp16 at load
  4fa4c7f 2026-09-03 BlivionIaG: debug(exl3): log suh/svh norm in marker debug to diagnose mul1 scale
  da36e0f 2026-09-03 BlivionIaG: debug(exl3): log CG-PATH output norm for layer 17 to diagnose mul1 inf
  6f66160 2026-09-03 BlivionIaG: fix(exl3): restore defensive activation buffer checks removed by sync commits
  8ea9db8 2026-09-03 BlivionIaG: fix(exl3): use torch.zeros for lm_head dequant output buffer
  6acee1b 2026-09-04 BlivionIaG: perf(exl3): extend grain-based v2 dispatch to full decode batches (M<=8)
  89a1513 2026-09-04 BlivionIaG: perf(exl3): grain-based M<=2 decode GEMM kernel (6.7x decode speedup)
  2240c99 2026-09-04 BlivionIaG: fix(exl3): runtime M-dispatch via custom op to stop decode regression
  7802315 2026-09-04 BlivionIaG: fix(exl3): load-time scratch pool for the decode prefill path
  62694f2 2026-09-04 BlivionIaG: feat(exl3): decode-trellis prefill path (7.5x faster prefill GEMM)
  4686661 2026-09-05 BlivionIaG: docs(recipe): linear-attention state breakdown (GDN per-request cost)
  566f529 2026-09-05 BlivionIaG: docs(recipe): Steam Deck VRAM fitment matrix + throughput expectation
  8f4d302 2026-09-05 BlivionIaG: docs(recipe): Ornith-EXL3-gfx1030 setup recipe (perf, VRAM, Steam Deck fitment)
  c7053d4 2026-09-05 BlivionIaG: perf(exl3): drop CG-PATH duplicate trellis/svh buffers (-2.45 GiB, 0% perf)
  c3e5258 2026-09-05 BlivionIaG: feat(exl3): VLLM_EXL3_MEMORY_MODE env var (full/packed), stub dispatcher
  3d6df9e 2026-09-18 BlivionIaG: fix(rocm): extend exl3_dot2 RDNA arch guard to all RDNA arches in docker build
  bc5fbee 2026-10-01 BlivionIaG: feat(rdna2): one serve entry point + per-config recipes
  b13effc 2026-10-01 BlivionIaG: fix(exl3): serve the quantized MTP draft on gfx1030 (MTP=2)
  6f1a712 2026-10-01 BlivionIaG: feat(rdna2): named TunableOp profiles + dedup/user-agnostic launchers
  d0df2f2 2026-10-01 BlivionIaG: feat(rdna2): EXL3 27B expansion of shared TunableOp set (760 -> 783)
  2818987 2026-10-01 BlivionIaG: feat(rdna2): EXL3 27B TunableOp capture/tune/curate/freeze tooling
  74315f4 2026-10-01 BlivionIaG: fix(rocm): gemma_rms_norm fake kernel contiguous output + exl3 27B mul1 TP=4 validation
  30b0bd4 2026-10-01 BlivionIaG: test(exl3): backport loader unit tests + mul1 tile reformat
  c7bc938 2026-10-01 BlivionIaG: feat(exl3): backport EXL3 n-gram table gather (PLE)
  a50520a 2026-10-01 BlivionIaG: feat(exl3): backport EXL3 MoE method for RoutedExperts
  aaa374a 2026-10-01 BlivionIaG: feat(exl3): backport v0.30 loader rewrite (mul1/mcg codebook, TP slicing, suh grouping)
  7be99d8 2026-10-01 BlivionIaG: feat(rdna): add exl3_project_rdna2 fused decode projection op
  a5059cf 2026-10-01 BlivionIaG: fix(rdna2): bound exl3 MoE trellis staging to the block's n-range
  0830a85 2026-10-01 BlivionIaG: fix(exl3): fold the 6-bit lm_head through the trellis decode
  2fb1d41 2026-10-01 BlivionIaG: port(exl3): mul1 decode and K=1..8 trellis kernels from v0.30
@BlivionIaG @GeorgeMurAlkh
feat(rdna2): FA-RDNA2 attention, GDN and Mamba paths
e8d1701 
FA-RDNA2 / RDNA_ATTN backend, GDN metadata (PIECEWISE state arenas),
Qwen GDN custom op, Mamba helpers. GDN cross-group metadata reuse
(vllm-project#58762) is disabled while static state buffers are active.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 14 file(s):
  3c85702 2026-07-15 BlivionIaG: feat(rocm): Wire FA-RDNA2 dispatch into RocmAttentionImpl
  7c3fec8 2026-07-31 BlivionIaG: feat(rocm): enable gfx1030 (RDNA2) as a recognized RDNA arch
  40568c8 2026-08-09 BlivionIaG: feat(rocm): stub rocm_rdna2_mla_sparse.py for gfx1030 sparse MLA decode (fallthrough to Triton)
  ea4e61a 2026-08-23 BlivionIaG: feat: hand-port GDN chunked prefill to HIP RDNA2 (5 kernels)
  5552701 2026-08-23 BlivionIaG: feat: hand-port GDN packed decode to HIP RDNA2
  126ecfa 2026-08-31 BlivionIaG: rdna: exl3 6bpw lm_head NotImplementedError + RDNA warmup guard
  e93cca1 2026-09-02 BlivionIaG: rdna: backport FA-RDNA2 test files (writer_layout + shape_sweep)
  5206533 2026-09-02 BlivionIaG: rdna: port K192 RDNA2 kernel work (W8A16 + W8A16-FP8 + sparse MLA)
  3859586 2026-09-03 BlivionIaG: fix(attention): add missing sliding_window arg to fa_rdna2 splitk wrapper
  f936195 2026-09-03 BlivionIaG: fix(gdn): skip host-sync NaN guard during cudagraph capture
  e85909f 2026-09-03 BlivionIaG: feat(attention): add standalone RDNA_ATTN backend with proven FA-RDNA2 kernels
  30d9f71 2026-09-03 BlivionIaG: debug(gdn): log out_buf nan check after kernel call
  e69cad4 2026-09-03 BlivionIaG: fix(gdn): simplify ssm_state zeroing condition (nan OR all-zeros)
  c48c86a 2026-09-03 BlivionIaG: fix(gdn): zero ssm_state before RDNA2 kernel call to avoid NaN propagation
  abc77aa 2026-09-04 BlivionIaG: perf(attn): fa_rdna2 decode kv_splits 8 -> 16
  e98975b 2026-09-04 BlivionIaG: fix(gdn): one-shot ssm_state page-commit scan + decode kernel toggle
  aac1fcd 2026-09-04 BlivionIaG: fix(gdn): chunk-local indexing bug in delta_h prefill kernel
  5ce7d86 2026-09-04 BlivionIaG: fix(gdn): default GDN HIP prefill chain to off (chunk-boundary corruption)
  1eb5569 2026-09-06 BlivionIaG: Revert "fix(rocm): auto-skip RDNA_ATTN MTP-verify gate when speculative decoding off"
  3a5e20b 2026-09-06 BlivionIaG: fix(rocm): auto-skip RDNA_ATTN MTP-verify gate when speculative decoding off
  cac4392 2026-09-06 BlivionIaG: fix(rocm): make RDNA_ATTN MTP-verify gate bypassable for large cudagraph capture
  f46bde9 2026-09-06 BlivionIaG: fix(gdn): drop ssm_state NaN check from VLLM_GDN_DBG debug print
  cafe95e 2026-09-08 BlivionIaG: fix(rocm): FPP13 checkpoint — FULL_AND_PIECEWISE greedy is coherent
  ba4c50d 2026-09-08 BlivionIaG: feat(gdn): VLLM_LOG_GDN_PTRS env-gated pointer instrumentation
  9e7184d 2026-09-08 BlivionIaG: fix(rocm,rdna-attn): disable spec gate by default, opt-in via env var
  f9c956a 2026-09-08 BlivionIaG: feat(rocm,fa-rdna2): prefer torch.ops._rocm_C over load_inline
  3787112 2026-09-08 BlivionIaG: Reapply "feat(rocm,rdna-attn): eager_break_during_capture on do_kv_cache_update"
  2873482 2026-09-08 BlivionIaG: Revert "feat(rocm,rdna-attn): eager_break_during_capture on do_kv_cache_update"
  13a5429 2026-09-08 BlivionIaG: feat(rocm,rdna-attn): eager_break_during_capture on do_kv_cache_update
  5eb84b4 2026-09-08 BlivionIaG: feat(rocm): AWQ prefill kernel + causal_conv1d HIP + platform gates
  631aa35 2026-09-08 BlivionIaG: feat(rocm,gdn): eager_break_during_capture for Qwen3.5/3.8 GDN + architecture whitelist
  7307db0 2026-09-08 BlivionIaG: fix(rocm,gdn-prefill): torch.zeros for cudagraph-safe scratch allocations
  2b2e1fb 2026-09-08 BlivionIaG: Revert "fix(rocm): revert rdna_attn capture override + kv_splits=16 to 0.27.1 form"
  8519ce7 2026-09-08 BlivionIaG: fix(rocm): revert rdna_attn capture override + kv_splits=16 to 0.27.1 form
  74f47b6 2026-09-10 BlivionIaG: fix(gdn): assign state arenas directly instead of register_buffer
  6ffcdec 2026-09-10 BlivionIaG: fix(gdn): pre-capture arena alloc and static-only gather slots
  dd4fb0b 2026-09-10 BlivionIaG: fix(gdn): gather/scatter by decode BS; lazy prefill backend
  369773e 2026-09-10 BlivionIaG: fix(gdn): freeze conv/ssm state arenas for piecewise cudagraph
  13a3daf 2026-09-11 BlivionIaG: cleanup(gqa): drop true-GQA variant, keep subgroup + off
  ecfec4e 2026-09-11 BlivionIaG: perf: GQA multi-head FA-RDNA2 prefill kernels for gfx1030
  4864d72 2026-09-12 BlivionIaG: debug(rocm): extend GDN probe with buffer identity + L6 arg capture
  b205fc0 2026-09-12 BlivionIaG: debug(rocm): env-gated forensics probes for TP=4 decode NaN
  854dd56 2026-09-12 BlivionIaG: fix(rocm): sanitize GDN state caches under cudagraph capture
  82b6f18 2026-09-13 BlivionIaG: fix(rocm): enable HIP causal_conv1d_fwd_rdna2 (out-stride + null-block)
  a447b23 2026-09-13 BlivionIaG: fix(rocm): use HIP reshape_and_cache_flash_rdna2 for the non-native KV write
  7259f0f 2026-09-14 BlivionIaG: fix(rocm,qwen4_exp): make the Flash-Next PLE mamba path load on gfx1030
  b2584f0 2026-09-15 George Muravei-Alkhavoi: fix(rocm): restore V620 CPU PLE and MTP startup compatibility
  143e2bf 2026-09-15 BlivionIaG: fix(rocm): zeros_like for GDN prefill output fallback (RDNA2 page-commit)
  02adbfd 2026-09-15 BlivionIaG: feat(rocm): fp16 state support in gdn_decode_rdna2
  ce3c939 2026-09-15 BlivionIaG: feat(rocm): select native RDNA2 HIP GDN prefill chain on gfx10x
  faf1040 2026-09-16 BlivionIaG: debug(rocm): gated NaN/copy instrumentation for FULL-graph corruption
  388a61b 2026-09-17 BlivionIaG: fix(rocm,gdn): remove the one-shot ssm/conv state zero_() wipes that corrupted the first request
  cd1231f 2026-09-17 BlivionIaG: perf(rocm,gdn): default the GDN prefill to Triton/FLA
  9dce178 2026-09-27 BlivionIaG: perf(rocm): FA-RDNA2 decode kernel for spec-decode verify and mixed batches
  690cbf0 2026-09-27 BlivionIaG: perf(rocm): FA-RDNA2 causal tile skipping, in-place output, parallel GQA softmax
  97fbae5 2026-09-27 BlivionIaG: [ROCm][RDNA] Load and call hippihx V1 from RDNA_ATTN decode (opt-in)
  e0112c5 2026-09-28 BlivionIaG: fix(rocm): gate split decode above the measured row threshold
  83e6af8 2026-09-28 BlivionIaG: perf(rdna2): fastest HIP attention path - FA verify-decode + verify capture; fa default
  b13effc 2026-10-01 BlivionIaG: fix(exl3): serve the quantized MTP draft on gfx1030 (MTP=2)
  284fbad 2026-10-02 BlivionIaG: [FA-RDNA2][rdna2] gqa prefill kernel templated on HEAD_DIM (128/256)

Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
@leapdragon @eastwood-c @BlivionIaG
feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-und…
6a314f5 
…er-PP

FA-RDNA2 / RDNA_ATTN backend, GDN metadata (PIECEWISE state arenas),
Qwen GDN custom op, Mamba helpers. GDN cross-group metadata reuse
(vllm-project#58762) is disabled while static state buffers are active.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 2 file(s):
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP

Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
Co-authored-by: BlivionIaG <kev29lt@gmail.com>
@Karl0007 @BlivionIaG
[Bugfix] Index mamba spec-decode block tables by request slot, not ba…
ba94c4f 
…tch row

FA-RDNA2 / RDNA_ATTN backend, GDN metadata (PIECEWISE state arenas),
Qwen GDN custom op, Mamba helpers. GDN cross-group metadata reuse
(vllm-project#58762) is disabled while static state buffers are active.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 1 file(s):
  741e5bc 2026-09-18 Karl0007: [Bugfix] Index mamba spec-decode block tables by request slot, not batch row

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
@BlivionIaG @leapdragon
feat(rocm): rdna_ar one-shot all-reduce for gfx10x PCIe
01e259c 
rdna_ar one-shot all-reduce (VLLM_RDNA_AR=1) placed first in the
CudaCommunicator all-reduce chain, plus custom-AR topology helpers.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 9 file(s):
  fb1687b 2026-09-03 BlivionIaG: feat(rocm): VLLM_FORCE_CUSTOM_ALL_REDUCE for PCIe-only RDNA
  cafe95e 2026-09-08 BlivionIaG: fix(rocm): FPP13 checkpoint — FULL_AND_PIECEWISE greedy is coherent
  a406064 2026-09-08 BlivionIaG: [review-only] cherry-pick leapdragon rdna_ar fabric knobs onto rdna_extras (#1)
  c091c42 2026-09-10 BlivionIaG: wip: checkpoint-validated state + chase csrc + debug instruments + row-zero fix
  e0a7b16 2026-09-12 BlivionIaG: fix(rdna_ar): make one-shot AR boot-robust, keep it opt-in
  965c3b9 2026-09-13 BlivionIaG: feat(rdna): add PIX topology helpers without changing AR policy
  849292e 2026-09-16 BlivionIaG: fix(rocm): make custom all-reduce correct under breakable cudagraphs
  c350fa2 2026-09-16 BlivionIaG: merge(rdna2): port PR #11 wvSplitK ownership/decode + rdna_ar device fix
  3b59ee1 2026-09-18 BlivionIaG: feat(rocm): opt-in leap T44b rdna_ar (VRAM flags + wedge check) (#13)
  295a842 2026-09-24 BlivionIaG: feat(rocm): make rdna_ar cover the custom-allreduce cases
  22d6e34 2026-09-27 BlivionIaG: feat(rdna2): two-shot rdna_ar (PR #22) + wide gate default, validated on 4x V620
  101a16c 2026-09-28 BlivionIaG: fix(rdna2): disable the racy two-shot rdna_ar by default; add batching knobs

Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
@BlivionIaG @GeorgeMurAlkh
@leapdragon @Karl0007 @wsantos @eastwood-c
fix(rocm): engine, cudagraph, scheduler and KV-cache fixes for gfx1030
4b8444c 
V1/V2 runner, cudagraph and scheduler fixes for gfx1030: ROCm FULL replay
inside the forward context / FULL-as-piecewise, capture-output growth,
rdna2 persist capture hooks, no UVA pointers in captured graphs, mrope
contiguous positions, TP>2 Mamba aligned-split gate (ROCm only),
adaptive prefill cap feeding upstream's per-step threshold, replay
boundaries as split stops, prefix-cache debug traces.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 18 file(s):
  07f1abf 2026-08-24 BlivionIaG: fix(distributed): allow TP comm ops in torch.compile graph capture
  f955842 2026-09-05 BlivionIaG: feat(debug): add per-step phase timing in v2 GPU model runner
  cafe95e 2026-09-08 BlivionIaG: fix(rocm): FPP13 checkpoint — FULL_AND_PIECEWISE greedy is coherent
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP
  b7c77f0 2026-09-11 BlivionIaG: perf: gate per-replay NaN input scan behind VLLM_CG_NAN_INPUT_CHECK
  b205fc0 2026-09-12 BlivionIaG: debug(rocm): env-gated forensics probes for TP=4 decode NaN
  f776159 2026-09-12 BlivionIaG: fix(rocm): wire rdna2 persist capture hooks around cudagraph capture
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  7259f0f 2026-09-14 BlivionIaG: fix(rocm,qwen4_exp): make the Flash-Next PLE mamba path load on gfx1030
  b2584f0 2026-09-15 George Muravei-Alkhavoi: fix(rocm): restore V620 CPU PLE and MTP startup compatibility
  3e1a0e1 2026-09-15 BlivionIaG: debug(rocm): gate mamba align-copy tracing behind VLLM_MAMBA_COPY_DEBUG
  8dc656b 2026-09-15 BlivionIaG: fix(rocm): cudagraph replay GPU page fault on UVA staging buffers
  849292e 2026-09-16 BlivionIaG: fix(rocm): make custom all-reduce correct under breakable cudagraphs
  faf1040 2026-09-16 BlivionIaG: debug(rocm): gated NaN/copy instrumentation for FULL-graph corruption
  ed94e3f 2026-09-18 BlivionIaG: fix(rocm,moe,spec): Qwen4Exp MTP path prerequisites from the MTP=2 investigation
  3b59ee1 2026-09-18 BlivionIaG: feat(rocm): opt-in leap T44b rdna_ar (VRAM flags + wedge check) (#13)
  741e5bc 2026-09-18 Karl0007: [Bugfix] Index mamba spec-decode block tables by request slot, not batch row
  fea652b 2026-09-22 George Muravei-Alkhavoi: Bound recurrent slot mappings during long V620 prefills
  6c26c78 2026-09-24 BlivionIaG: fix(rocm): keep FULL decode graphs alongside piecewise capture
  30632b2 2026-09-27 BlivionIaG: perf(spec-decode): capture the draft decode graph for max_num_reqs
  6ed3909 2026-09-28 BlivionIaG: feat(engine): prefill step cadence for non-DP engines
  e3bf3d3 2026-10-05 BlivionIaG: feat(sched): restore upstream mixed decode+prefill batching

Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
Co-authored-by: Karl0007 <44745209+Karl0007@users.noreply.github.com>
Co-authored-by: Waldecir Santos <waldecir@gmail.com>
Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
@GeorgeMurAlkh @BlivionIaG @wsantos
fix(rocm): port George Muravei-Alkhavoi's engine changes (10 commits)
f69aaf8 
V1/V2 runner, cudagraph and scheduler fixes for gfx1030: ROCm FULL replay
inside the forward context / FULL-as-piecewise, capture-output growth,
rdna2 persist capture hooks, no UVA pointers in captured graphs, mrope
contiguous positions, TP>2 Mamba aligned-split gate (ROCm only),
adaptive prefill cap feeding upstream's per-step threshold, replay
boundaries as split stops, prefix-cache debug traces.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 10 file(s):
  c091c42 2026-09-10 BlivionIaG: wip: checkpoint-validated state + chase csrc + debug instruments + row-zero fix
  b205fc0 2026-09-12 BlivionIaG: debug(rocm): env-gated forensics probes for TP=4 decode NaN
  f3f4ad0 2026-09-12 BlivionIaG: fix(rocm): gate mamba block-aligned prefill split to TP<=2
  e45dd5c 2026-09-17 BlivionIaG: fix(v1,core): let CircularBufferManager accept group-aligned prefix hits
  772dac4 2026-09-22 George Muravei-Alkhavoi: Port served V620 prefix checkpoint replay fixes
  fea652b 2026-09-22 George Muravei-Alkhavoi: Bound recurrent slot mappings during long V620 prefills
  b88acd9 2026-09-23 George Muravei-Alkhavoi: Revert "Instrument V620 scheduler block pressure"
  5030766 2026-09-23 George Muravei-Alkhavoi: Revert "Trace V620 cache group allocation"
  249a565 2026-09-23 George Muravei-Alkhavoi: Trace V620 cache group allocation
  80c9874 2026-09-23 George Muravei-Alkhavoi: Instrument V620 scheduler block pressure
  f614309 2026-09-26 George Muravei-Alkhavoi: test(kv-cache): correct scratch replay boundary
  f6e5259 2026-09-26 George Muravei-Alkhavoi: fix(kv-cache): skip scratch groups when truncating hits
  e690365 2026-09-26 George Muravei-Alkhavoi: test(kv-offload): align ROCm and MTP regressions
  5fa2d48 2026-09-26 George Muravei-Alkhavoi: fix(rocm): keep CPU KV offload buffers private
  1512a7c 2026-10-04 Waldecir Santos: feat(rdna2): adaptive prefill scheduling (dynamic chunk cap + step cadence)
  e3bf3d3 2026-10-05 BlivionIaG: feat(sched): restore upstream mixed decode+prefill batching

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
Co-authored-by: Waldecir Santos <waldecir@gmail.com>
@leapdragon @BlivionIaG
@GeorgeMurAlkh @eastwood-c
fix(rocm): port Aron Hsiao's engine changes (2 commits)
d5b3372 
V1/V2 runner, cudagraph and scheduler fixes for gfx1030: ROCm FULL replay
inside the forward context / FULL-as-piecewise, capture-output growth,
rdna2 persist capture hooks, no UVA pointers in captured graphs, mrope
contiguous positions, TP>2 Mamba aligned-split gate (ROCm only),
adaptive prefill cap feeding upstream's per-step threshold, replay
boundaries as split stops, prefix-cache debug traces.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 6 file(s):
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  4224ce2 2026-09-15 BlivionIaG: fix(rocm): cache mamba state copy funcs to fix per-step singular/plural bug
  8dc656b 2026-09-15 BlivionIaG: fix(rocm): cudagraph replay GPU page fault on UVA staging buffers
  1ff7359 2026-09-16 BlivionIaG: fix(rocm): make V1 FULL_AND_PIECEWISE + prefix caching correct on gfx1030
  4425834 2026-09-18 BlivionIaG: fix(rocm): make platform detection and worker init robust on RDNA consumer GPUs
  3b59ee1 2026-09-18 BlivionIaG: feat(rocm): opt-in leap T44b rdna_ar (VRAM flags + wedge check) (#13)
  6c26c78 2026-09-24 BlivionIaG: fix(rocm): keep FULL decode graphs alongside piecewise capture

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
@wsantos @BlivionIaG
feat(rdna2): adaptive prefill scheduling (dynamic chunk cap + step ca…
dd85ff2 
…dence)

V1/V2 runner, cudagraph and scheduler fixes for gfx1030: ROCm FULL replay
inside the forward context / FULL-as-piecewise, capture-output growth,
rdna2 persist capture hooks, no UVA pointers in captured graphs, mrope
contiguous positions, TP>2 Mamba aligned-split gate (ROCm only),
adaptive prefill cap feeding upstream's per-step threshold, replay
boundaries as split stops, prefix-cache debug traces.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 2 file(s):
  6c26c78 2026-09-24 BlivionIaG: fix(rocm): keep FULL decode graphs alongside piecewise capture
  1512a7c 2026-10-04 Waldecir Santos: feat(rdna2): adaptive prefill scheduling (dynamic chunk cap + step cadence)

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
@BlivionIaG
feat(rdna2): Qwen3.8 Flash-Next (Qwen4Exp) and DeepSeek-V4 on gfx1030
95aed84 
Qwen3.8 Flash-Next (Qwen4Exp) RDNA2 deltas on top of upstream's
qwen4_exp (FP16 QSA, HIP QSA/HC/PLE-conv opt-ins, eager break),
DeepSeek-V4 RDNA2 sparse-MLA decode routing, MTP local argmax,
unquantized-MTP detection.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 6 file(s):
  40568c8 2026-08-09 BlivionIaG: feat(rocm): stub rocm_rdna2_mla_sparse.py for gfx1030 sparse MLA decode (fallthrough to Triton)
  5eea5c9 2026-08-28 BlivionIaG: wip(rocm): exl3 3inst decode working end-to-end (Qwen3.5-0.8B)
  1eb11e1 2026-09-14 BlivionIaG: feat(rocm): opt-in HIP paths for Qwen4Exp / Qwen3.8-Flash-Next on gfx1030
  b26763e 2026-09-15 BlivionIaG: fix(rocm): ple_short_conv optional-tensor schema + on_gfx10x import path
  8960a3b 2026-09-16 BlivionIaG: fix(rocm): make the Qwen4Exp HC HIP chain compute correctly
  3092d63 2026-09-16 BlivionIaG: fix(rocm): new_zeros for hc_combine_norm/hc_combine output buffers
  c59a23f 2026-09-17 BlivionIaG: chore: recover remote-only work products into the repo
  50120e1 2026-09-17 BlivionIaG: fix(rocm,hc): make the HC HIP dispatcher's _contig() capture-safe
  d19b87e 2026-10-02 BlivionIaG: port: QSA RDNA2 HIP decode path (v0.30 fix files onto rdna_extras)
  c6d99ca 2026-10-06 BlivionIaG: fix(rdna2): HC HIP stream correctness + MoE prefill tile
@leapdragon @BlivionIaG
@GeorgeMurAlkh @eastwood-c
feat(rdna2): port Aron Hsiao's qwen4exp-dsv4 changes (2 commits)
92b2cb1 
Qwen3.8 Flash-Next (Qwen4Exp) RDNA2 deltas on top of upstream's
qwen4_exp (FP16 QSA, HIP QSA/HC/PLE-conv opt-ins, eager break),
DeepSeek-V4 RDNA2 sparse-MLA decode routing, MTP local argmax,
unquantized-MTP detection.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 12 file(s):
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  1eb11e1 2026-09-14 BlivionIaG: feat(rocm): opt-in HIP paths for Qwen4Exp / Qwen3.8-Flash-Next on gfx1030
  b2584f0 2026-09-15 George Muravei-Alkhavoi: fix(rocm): restore V620 CPU PLE and MTP startup compatibility
  8cf0ded 2026-09-15 BlivionIaG: fix(rocm,qwen4_exp): Qwen3.8-Flash-Next Triton path serves in eager mode (capture break point)
  fa1cf12 2026-09-15 BlivionIaG: fix(rocm,qwen4_exp): Qwen3.8-Flash-Next serves in eager mode (PR#5 cache-writer) + status
  1ff7359 2026-09-16 BlivionIaG: fix(rocm): make V1 FULL_AND_PIECEWISE + prefix caching correct on gfx1030
  faf1040 2026-09-16 BlivionIaG: debug(rocm): gated NaN/copy instrumentation for FULL-graph corruption
  bb4cb3b 2026-09-16 BlivionIaG: fix(rocm): make PleOffloadLayer forward an eager capture break point
  e7e6e35 2026-09-20 George Muravei-Alkhavoi: style(rocm): normalize modified QSA modules
  de6f1c9 2026-09-20 George Muravei-Alkhavoi: Carry QSA live context into side metadata
  9a91ff6 2026-09-20 George Muravei-Alkhavoi: Bound AMD QSA scoring to live prefill context
  dbd9994 2026-09-22 George Muravei-Alkhavoi: Fuse AMD PLE grouped normalization on V620
  3511821 2026-09-23 George Muravei-Alkhavoi: Revert "Instrument long-context QSA stages"
  2f93cd1 2026-09-23 George Muravei-Alkhavoi: Revert "Skip QSA timing sync during graph capture"
  17952ec 2026-09-23 George Muravei-Alkhavoi: Skip QSA timing sync during graph capture
  77823e1 2026-09-23 George Muravei-Alkhavoi: Instrument long-context QSA stages
  6c26c78 2026-09-24 BlivionIaG: fix(rocm): keep FULL decode graphs alongside piecewise capture
  700753d 2026-09-25 BlivionIaG: fix(rocm,mtp): unbreak MTP draft inference — amdsmi shutdown masking + local-argmax draft sampling
  d19b87e 2026-10-02 BlivionIaG: port: QSA RDNA2 HIP decode path (v0.30 fix files onto rdna_extras)

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
@GeorgeMurAlkh @BlivionIaG
Fuse AMD PLE grouped normalization on V620
83ef316 
Qwen3.8 Flash-Next (Qwen4Exp) RDNA2 deltas on top of upstream's
qwen4_exp (FP16 QSA, HIP QSA/HC/PLE-conv opt-ins, eager break),
DeepSeek-V4 RDNA2 sparse-MLA decode routing, MTP local argmax,
unquantized-MTP detection.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 1 file(s):
  dbd9994 2026-09-22 George Muravei-Alkhavoi: Fuse AMD PLE grouped normalization on V620
A17T - Simon Siebert @BlivionIaG
@leapdragon @eastwood-c
fix(models): detect unquantized MTP weights from checkpoint metadata
7306979 
Qwen3.8 Flash-Next (Qwen4Exp) RDNA2 deltas on top of upstream's
qwen4_exp (FP16 QSA, HIP QSA/HC/PLE-conv opt-ins, eager break),
DeepSeek-V4 RDNA2 sparse-MLA decode routing, MTP local argmax,
unquantized-MTP detection.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 1 file(s):
  22bb2e8 2026-09-11 Aron Hsiao: feat(rdna2): port recipe Triton LDS/softmax, skinny MoE GEMV, MTP-under-PP
  e86265f 2026-09-22 A17T - Simon Siebert: fix(models): detect unquantized MTP weights from checkpoint metadata
  b13effc 2026-10-01 BlivionIaG: fix(exl3): serve the quantized MTP draft on gfx1030 (MTP=2)

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
Co-authored-by: Chris Eastwood <chris.eastwood@pwn4g3.dev>
@BlivionIaG @leapdragon
feat(rocm): PLE CPU sidecar offload (kept, not wired on v0.31)
cd1c549 
PLE CPU sidecar offload (vllm/v1/ple_offload). Kept in tree but NOT
wired into Qwen4Exp on v0.31: upstream now offloads the PLE table via
EngramConfig.cpu_offload (pinned host UVA). VLLM_PLE_CPU_OFFLOAD=1 is
rejected with a clear error until this is decided.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 1 file(s):
  16695b2 2026-09-18 BlivionIaG: test(rocm): isolate V620 amdsmi and PLE fp8 CPU tests from HIP
  b8354ff 2026-09-18 BlivionIaG: feat(rocm): V620 MoE JSON, ROCR amdsmi index, PLE fp8 fused gather

Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
@leapdragon @BlivionIaG @GeorgeMurAlkh
feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
460da29 
PLE CPU sidecar offload (vllm/v1/ple_offload). Kept in tree but NOT
wired into Qwen4Exp on v0.31: upstream now offloads the PLE table via
EngramConfig.cpu_offload (pinned host UVA). VLLM_PLE_CPU_OFFLOAD=1 is
rejected with a clear error until this is decided.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 7 file(s):
  5765f57 2026-09-12 Aron Hsiao: feat(rdna2): port Qwen4Exp / Qwen3.8 Flash-Next for gfx1030
  b85ebaa 2026-09-13 BlivionIaG: fix(rocm,ple): guard missing ngram_embedding + refresh stale PLE routing test
  469c456 2026-09-14 BlivionIaG: fix(rocm,ple): resolve HIP through torch on gfx1030 (709) + overridable block size
  e7ce541 2026-09-14 BlivionIaG: fix(rocm,ple): make Flash-Next PLE offload share CUDA tensors on gfx1030
  b2584f0 2026-09-15 George Muravei-Alkhavoi: fix(rocm): restore V620 CPU PLE and MTP startup compatibility
  7bc1c36 2026-09-15 BlivionIaG: debug(rocm): capture-safe PLE output NaN scan (VLLM_PLE_NAN_DEBUG=1)
  bb4cb3b 2026-09-16 BlivionIaG: fix(rocm): make PleOffloadLayer forward an eager capture break point
  b8354ff 2026-09-18 BlivionIaG: feat(rocm): V620 MoE JSON, ROCR amdsmi index, PLE fp8 fused gather

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
@BlivionIaG @GeorgeMurAlkh
@wsantos @leapdragon
tools: RDNA launchers, recipes, probes and bench drivers
5ab3033 
Launchers (scripts/serve_rdna.sh, recipes), probes, bench drivers and
kernel tests/tools. Launchers stay in scripts/ (upstream moved its own
auxiliary files to tools/, vllm-project#58572).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 75 file(s):
  4c3d877 2026-09-06 BlivionIaG: feat(ops): add sync_remote_to_local.sh for validated-state capture
  c091c42 2026-09-10 BlivionIaG: wip: checkpoint-validated state + chase csrc + debug instruments + row-zero fix
  13a3daf 2026-09-11 BlivionIaG: cleanup(gqa): drop true-GQA variant, keep subgroup + off
  dadfeab 2026-09-11 BlivionIaG: chore: surface VLLM_FA_RDNA2_GQA_MODE in production launcher
  d74eb45 2026-09-11 BlivionIaG: chore: make AR flags overridable in serve_gfx1030_full.sh
  e0a7b16 2026-09-12 BlivionIaG: fix(rdna_ar): make one-shot AR boot-robust, keep it opt-in
  ca83d92 2026-09-13 BlivionIaG: feat(rdna2): import PR #5 V620 TunableOp tables + tuning tooling
  965c3b9 2026-09-13 BlivionIaG: feat(rdna): add PIX topology helpers without changing AR policy
  f86faad 2026-09-13 BlivionIaG: fix(rocm): enable RDNA2 W4A16 at TP=4 via breakable cudagraphs + PYNCCL
  469c456 2026-09-14 BlivionIaG: fix(rocm,ple): resolve HIP through torch on gfx1030 (709) + overridable block size
  e7ce541 2026-09-14 BlivionIaG: fix(rocm,ple): make Flash-Next PLE offload share CUDA tensors on gfx1030
  59237b3 2026-09-14 BlivionIaG: fix(rdna2): repair 0008 skinny default, gfx10 bf16 abort, warmup timeout
  f663686 2026-09-15 BlivionIaG: chore(rdna2): drop vestigial combo_kernels=false from serve config
  7e70e24 2026-09-16 BlivionIaG: fix(scripts): bench_report selects the newest result by mtime
  f5cbbdf 2026-09-16 BlivionIaG: feat(scripts): bench_report.py for PP/TG/TTFT reporting
  4d25a04 2026-09-16 BlivionIaG: update(rocm): default custom all-reduce on in the gfx1030 launcher
  1ff7359 2026-09-16 BlivionIaG: fix(rocm): make V1 FULL_AND_PIECEWISE + prefix caching correct on gfx1030
  657bdba 2026-09-17 BlivionIaG: perf(scripts): raise KV pool to 7 GiB and GPU memory utilization to 0.90
  0dd3811 2026-09-17 BlivionIaG: feat(scripts): Flash-Next launcher serves native 262k context with vision on
  aaa4775 2026-09-17 BlivionIaG: fix(scripts): use the official qwen3_coder tool parser for Flash-Next
  37345ee 2026-09-17 BlivionIaG: feat(scripts): enable tool calling + reasoning parse on the Flash-Next launcher
  ae5c7ed 2026-09-17 BlivionIaG: feat(scripts): serve Flash-Next under a stable served-model-name
  c59a23f 2026-09-17 BlivionIaG: chore: recover remote-only work products into the repo
  31003ff 2026-09-17 BlivionIaG: docs(scripts): document the validated vision-on stopgap in the Flash-Next launcher
  3bddd3c 2026-09-17 BlivionIaG: feat(scripts): add the validated Flash-Next production launcher (PIECEWISE, cap 6)
  b78006a 2026-09-17 BlivionIaG: fix(scripts): cap gfx1030 production serve at max-num-seqs 6, batched-tokens 2048
  609c9c0 2026-09-18 BlivionIaG: feat(scripts): Flash-Next production on FULL_AND_PIECEWISE
  16695b2 2026-09-18 BlivionIaG: test(rocm): isolate V620 amdsmi and PLE fp8 CPU tests from HIP
  b8354ff 2026-09-18 BlivionIaG: feat(rocm): V620 MoE JSON, ROCR amdsmi index, PLE fp8 fused gather
  9c60294 2026-09-18 BlivionIaG: feat(scripts): enable --enable-prompt-tokens-details on Flash-Next
  3b59ee1 2026-09-18 BlivionIaG: feat(rocm): opt-in leap T44b rdna_ar (VRAM flags + wedge check) (#13)
  26d6c4b 2026-09-23 George Muravei-Alkhavoi: Reduce V620 vision pixel cap
  295a842 2026-09-24 BlivionIaG: feat(rocm): make rdna_ar cover the custom-allreduce cases
  22d6e34 2026-09-27 BlivionIaG: feat(rdna2): two-shot rdna_ar (PR #22) + wide gate default, validated on 4x V620
  7010756 2026-09-27 BlivionIaG: feat(rdna2): MTP=0/2 Flash-Next launcher + lookup-only TunableOp env helper
  b95036b 2026-09-28 BlivionIaG: feat(rocm): 27B dense launcher + keep the W4A8 attempt out of dynamo's graph
  6ed3909 2026-09-28 BlivionIaG: feat(engine): prefill step cadence for non-DP engines
  ae5bedf 2026-09-28 BlivionIaG: fix(rdna2): production launcher - pin its own tree + FA defaults
  83e6af8 2026-09-28 BlivionIaG: perf(rdna2): fastest HIP attention path - FA verify-decode + verify capture; fa default
  56744dd 2026-09-28 BlivionIaG: fix(rdna2): launcher safety+batching consistency (scoped cleanup, one-shot default, knobs)
  330b42a 2026-09-28 BlivionIaG: perf(rdna2): default MAXBAT=2048 (best measured prefill/decode fairness point)
  4ca5aaa 2026-09-28 BlivionIaG: perf(rdna2): default long-prefill chunking to 1024 for decode fairness
  101a16c 2026-09-28 BlivionIaG: fix(rdna2): disable the racy two-shot rdna_ar by default; add batching knobs
  2204957 2026-09-29 BlivionIaG: perf(rdna2): ship curated FP16 TunableOp rows (719/rank) for Flash-Next
  27437bf 2026-09-29 BlivionIaG: chore(rdna2): standardize TunableOp storage on shared build-keyed rows
  c9122de 2026-09-29 BlivionIaG: feat(rocm): fused MoE W4A8 sdot4 for gfx1030 (Flash-Next reachable)
  3050090 2026-09-29 BlivionIaG: fix(rdna2): preserve pre-W4A8 W4A16 split-K choice; W4A8 defaults to 0
  a641d65 2026-09-29 BlivionIaG: fix(rdna2): K_STEP-align prefill split-K (W4A8 fallback NaN) + W4A8 selector guard
  8de5025 2026-09-29 BlivionIaG: feat(rdna2): RDNA_AR knob + fresh-cache 27B W4A16 AR-off driver
  45e218e 2026-09-29 BlivionIaG: feat(rdna2): env-gated W4A8 shape diagnostic + 27B dense A/B tooling
  3418e49 2026-09-29 BlivionIaG: tools(rdna2_028): add m27b matrix runner (16k/1k, 1k/512 at c=1, c=8)
  adcd0f2 2026-09-30 BlivionIaG: bench(explore): PR30 harness fixes + V620 results (H1 holds, H3 fails)
  d9c927c 2026-09-30 BlivionIaG: perf(rdna2): freeze 760-row TunableOp set from full-matrix shape census (719 -> 760)
  ee4f379 2026-09-30 BlivionIaG: perf(rdna2): lookup-hit proof + curate evidence for the 754-row set
  1707cc0 2026-09-30 BlivionIaG: perf(rdna2): expand TunableOp rows by 35 (719 -> 754) for w4a16-fa-m0 capture arm
  bc5fbee 2026-10-01 BlivionIaG: feat(rdna2): one serve entry point + per-config recipes
  6f1a712 2026-10-01 BlivionIaG: feat(rdna2): named TunableOp profiles + dedup/user-agnostic launchers
  ac9dfd8 2026-10-01 BlivionIaG: feat(rdna2): unify TunableOp rows + wire all launchers through the helper
  b41f39e 2026-10-01 BlivionIaG: bench(rdna2): Flash-Next fp32-accum matrix + CAS A/B + W4A8 arm
  ab5ccf3 2026-10-02 BlivionIaG: docs(rdna): sweep not-FA/not-HIP gaps — verdicts + PR#33 recovery
  e3bf3d3 2026-10-05 BlivionIaG: feat(sched): restore upstream mixed decode+prefill batching

Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
Co-authored-by: Waldecir Santos <waldecir@gmail.com>
Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
@GeorgeMurAlkh @BlivionIaG
tools: port George Muravei-Alkhavoi's tools changes (22 commits)
84fe940 
Launchers (scripts/serve_rdna.sh, recipes), probes, bench drivers and
kernel tests/tools. Launchers stay in scripts/ (upstream moved its own
auxiliary files to tools/, vllm-project#58572).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 5 file(s):
  93265f9 2026-09-22 George Muravei-Alkhavoi: Document Git-only V620 fast service reproduction
  0821be8 2026-09-22 George Muravei-Alkhavoi: Port opt-in served resident skinny MoE decode
  1e68a8b 2026-09-22 George Muravei-Alkhavoi: Version served V620 MoE tuning and launch settings
  573e4c8 2026-09-22 George Muravei-Alkhavoi: Make resident MoE an explicit benchmark launch option
  9a634c7 2026-09-22 George Muravei-Alkhavoi: Port resident W4A16 MoE and FP16 shared experts to current base
  9daf24d 2026-09-22 George Muravei-Alkhavoi: Add reproducible V620 base benchmark launcher
  26d6c4b 2026-09-23 George Muravei-Alkhavoi: Reduce V620 vision pixel cap
  b88acd9 2026-09-23 George Muravei-Alkhavoi: Revert "Instrument V620 scheduler block pressure"
  80c9874 2026-09-23 George Muravei-Alkhavoi: Instrument V620 scheduler block pressure
  c9ed993 2026-09-23 George Muravei-Alkhavoi: Test warmed 3.75 GiB V620 cache
  945bba2 2026-09-23 George Muravei-Alkhavoi: Restore measured V620 service memory settings
  3511821 2026-09-23 George Muravei-Alkhavoi: Revert "Instrument long-context QSA stages"
  77823e1 2026-09-23 George Muravei-Alkhavoi: Instrument long-context QSA stages
  71ec2ce 2026-09-23 George Muravei-Alkhavoi: Restore fused shared expert in Git service
  b00ae65 2026-09-26 George Muravei-Alkhavoi: test(rdna2): benchmark multimodal RAM KV reload
  cdf8d05 2026-09-26 George Muravei-Alkhavoi: test(rdna2): benchmark RAM KV cache round trip
  32d2273 2026-09-26 George Muravei-Alkhavoi: chore(rdna2): allow external runtime cache directory
  af47324 2026-09-26 George Muravei-Alkhavoi: feat(rdna2): parameterize RAM KV offload size
  f22fee0 2026-09-27 George Muravei-Alkhavoi: perf(rocm): opt in to HIP MoE for concurrent MTP decode
  28046ee 2026-09-27 George Muravei-Alkhavoi: deploy: select Swift 1.5 AutoRound for the default service
  5371c10 2026-09-27 George Muravei-Alkhavoi: deploy: make 64 GiB RAM KV offload the V620 default
  a146fcf 2026-09-27 George Muravei-Alkhavoi: tools: allow skipping multimodal startup profiling
  27437bf 2026-09-29 BlivionIaG: chore(rdna2): standardize TunableOp storage on shared build-keyed rows

Co-authored-by: BlivionIaG <kev29lt@gmail.com>
@wsantos @BlivionIaG
tools: port Waldecir Santos's tools changes (2 commits)
ed255c9 
Launchers (scripts/serve_rdna.sh, recipes), probes, bench drivers and
kernel tests/tools. Launchers stay in scripts/ (upstream moved its own
auxiliary files to tools/, vllm-project#58572).

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 2 file(s):
  fc04460 2026-09-23 Waldecir Santos: feat(rdna2): V620 PIECEWISE launcher and measured report
  1512a7c 2026-10-04 Waldecir Santos: feat(rdna2): adaptive prefill scheduling (dynamic chunk cap + step cadence)
@BlivionIaG @GeorgeMurAlkh
docs: RDNA2 docs, bench artifacts and TunableOp rows
231b8df 
RDNA2 docs, bench_results artifacts, TunableOp rows and repo notes.

Ported from opengfx1030/vllm-rdna rdna_extras@c6d99ca305 onto upstream v0.31.0.
Original commits touching these 366 file(s):
  7c3fec8 2026-07-31 BlivionIaG: feat(rocm): enable gfx1030 (RDNA2) as a recognized RDNA arch
  c7053d4 2026-09-05 BlivionIaG: perf(exl3): drop CG-PATH duplicate trellis/svh buffers (-2.45 GiB, 0% perf)
  d71721c 2026-09-08 BlivionIaG: docs(profiling): cudagraph dispatcher routing finding
  598a3f8 2026-09-08 BlivionIaG: docs(profiling): forcing GDN decode core eager does not fix it
  34e123f 2026-09-08 BlivionIaG: docs(profiling): max-num-seqs 1 rules out state-index handling
  9ac93df 2026-09-08 BlivionIaG: docs(profiling): static buffers in piecewise cause the non-determinism
  677f48e 2026-09-08 BlivionIaG: docs(profiling): record disproven non_blocking race + config matrix
  c797740 2026-09-08 BlivionIaG: docs(profiling): narrow cudagraph garbage to GDN hybrid state path
  9e7184d 2026-09-08 BlivionIaG: fix(rocm,rdna-attn): disable spec gate by default, opt-in via env var
  6c5ff94 2026-09-10 BlivionIaG: docs: journal resolution - GDN state arenas fix 16k c=8
  aaaae85 2026-09-11 BlivionIaG: perf: route W4A16 AWQ prefill through GPTQ ConfigA kernel
  ca83d92 2026-09-13 BlivionIaG: feat(rdna2): import PR #5 V620 TunableOp tables + tuning tooling
  312c7bb 2026-09-14 BlivionIaG: docs(rdna2): both QSA kernels pass standalone; abort is in the QSA forward integration
  f543afb 2026-09-14 BlivionIaG: docs(rdna2): QSA sparse-attention kernel passes standalone; abort is indexer/forward integration
  8204653 2026-09-14 BlivionIaG: docs(rdna2): record Flash-Next QSA/attention abort diagnostic in the HIP tracker
  5c3c0c6 2026-09-14 BlivionIaG: docs(rdna2): 27B AWQ bench matrix (1k/512 + 16k/1k at c=1/4/8)
  469c456 2026-09-14 BlivionIaG: fix(rocm,ple): resolve HIP through torch on gfx1030 (709) + overridable block size
  1eb11e1 2026-09-14 BlivionIaG: feat(rocm): opt-in HIP paths for Qwen4Exp / Qwen3.8-Flash-Next on gfx1030
  e7ce541 2026-09-14 BlivionIaG: fix(rocm,ple): make Flash-Next PLE offload share CUDA tensors on gfx1030
  fcea021 2026-09-15 BlivionIaG: docs(rdna2): Flash-Next prefix caching needs PIECEWISE, not FULL_AND_PIECEWISE
  e285a2a 2026-09-15 BlivionIaG: docs(rdna2): Flash-Next prefix-caching corruption (mamba align mode)
  7a1535e 2026-09-15 BlivionIaG: docs(rdna2): GDN decode investigation - launch-bound, not a bottleneck
  391fb64 2026-09-15 BlivionIaG: docs(rdna2): HIP GDN prefill wired in, 16k c=8 no regression
  ecf21df 2026-09-15 BlivionIaG: docs(rdna2): Flash-Next RDNA2 W4A16 wins all 6 cells vs 27B baseline
  34d7268 2026-09-15 BlivionIaG: docs(rdna2): Flash-Next full bench matrix vs 27B baseline
  ffd88e7 2026-09-15 BlivionIaG: docs(rdna2): Flash-Next cudagraph fix verified + 27B comparison
  93d07a0 2026-09-15 BlivionIaG: docs(rdna2): Flash-Next cudagraph replay = GPU page fault (stale host buffer)
  242a76b 2026-09-15 BlivionIaG: docs(rdna2): Flash-Next eager baseline + 16k/capture remaining
  fa1cf12 2026-09-15 BlivionIaG: fix(rocm,qwen4_exp): Qwen3.8-Flash-Next serves in eager mode (PR#5 cache-writer) + status
  f273a5c 2026-09-15 BlivionIaG: docs(rdna2): QSA sparse attention passes with the real transposed+split cache layout too; both remaining QSA suspects are exonerated
  a3946a4 2026-09-15 BlivionIaG: docs(rdna2): all 3 QSA-adjacent kernels pass standalone; abort is warmup-forward integration
  c350fa2 2026-09-16 BlivionIaG: merge(rdna2): port PR #11 wvSplitK ownership/decode + rdna_ar device fix
  c59a23f 2026-09-17 BlivionIaG: chore: recover remote-only work products into the repo
  295a842 2026-09-24 BlivionIaG: feat(rocm): make rdna_ar cover the custom-allreduce cases
  6c26c78 2026-09-24 BlivionIaG: fix(rocm): keep FULL decode graphs alongside piecewise capture
  bcdaadd 2026-09-27 BlivionIaG: docs(rdna2): operational note for the two-shot rdna_ar all-reduce
  0b85857 2026-09-27 BlivionIaG: docs(rdna2): note the validated rdna_ar gate (20480) and the disabled-backend symptom
  48b61b9 2026-09-27 BlivionIaG: docs(rdna2): document the fastest validated MTP path for reproduction
  645435b 2026-09-27 BlivionIaG: tunableop: share MTP FP16 tuning rows for the 0.28.0 stack (rocblas-f30bb442e9b5)
  4d94aaf 2026-09-28 BlivionIaG: feat(rocm): wire W4A8 (int4 x int8 sdot4) as an opt-in dense W4A16 prefill replacement
  101a16c 2026-09-28 BlivionIaG: fix(rdna2): disable the racy two-shot rdna_ar by default; add batching knobs
  2204957 2026-09-29 BlivionIaG: perf(rdna2): ship curated FP16 TunableOp rows (719/rank) for Flash-Next
  27437bf 2026-09-29 BlivionIaG: chore(rdna2): standardize TunableOp storage on shared build-keyed rows
  c9122de 2026-09-29 BlivionIaG: feat(rocm): fused MoE W4A8 sdot4 for gfx1030 (Flash-Next reachable)
  3050090 2026-09-29 BlivionIaG: fix(rdna2): preserve pre-W4A8 W4A16 split-K choice; W4A8 defaults to 0
  b1fcdcc 2026-09-29 BlivionIaG: bench(rdna2): preserve W4A8 split-K K_STEP-alignment fix evidence
  2fb0204 2026-09-29 BlivionIaG: bench: record TP=4 W4A16 FULL-HIP RDNA_AR=0 bring-up (host reset, PCI SERR e68)
  8b11337 2026-09-29 BlivionIaG: bench: record w4a8-wiring matrix results (MTP=0/2, 16k/1k, 1k/512, c=1/8)
  2138382 2026-09-29 BlivionIaG: feat(rocm): proper vLLM-idiomatic W4A8 prefill integration
  adcd0f2 2026-09-30 BlivionIaG: bench(explore): PR30 harness fixes + V620 results (H1 holds, H3 fails)
  4b37b52 2026-09-30 BlivionIaG: docs(explore): RDNA2 W4A16 M-dispatch under torch.compile
  bfff202 2026-09-30 BlivionIaG: bench(rdna2): validate the frozen 760-row TunableOp set across the serving matrix
  d9c927c 2026-09-30 BlivionIaG: perf(rdna2): freeze 760-row TunableOp set from full-matrix shape census (719 -> 760)
  6d0e486 2026-09-30 BlivionIaG: bench(rdna2): archive STEP 2 MTP=2 baseline (Flash-Next, TP=4, F&P, FA-RDNA2, W4A16, frozen rows)
  ee4f379 2026-09-30 BlivionIaG: perf(rdna2): lookup-hit proof + curate evidence for the 754-row set
  6f1a712 2026-10-01 BlivionIaG: feat(rdna2): named TunableOp profiles + dedup/user-agnostic launchers
  ac9dfd8 2026-10-01 BlivionIaG: feat(rdna2): unify TunableOp rows + wire all launchers through the helper
  9cad879 2026-10-01 BlivionIaG: [rdna2][MoE-epilogue] default back to packed-fp16 CAS; fp32 scratch opt-in
  b41f39e 2026-10-01 BlivionIaG: bench(rdna2): Flash-Next fp32-accum matrix + CAS A/B + W4A8 arm
  ab5ccf3 2026-10-02 BlivionIaG: docs(rdna): sweep not-FA/not-HIP gaps — verdicts + PR#33 recovery
  bb40498 2026-10-02 BlivionIaG: docs(bench): use clean GPU-5 kernel A/B in PR#34 summary
  ef8f0a2 2026-10-02 BlivionIaG: bench(rocm): PR#34 FA-RDNA2 D=128 prefill validation on gfx1030

Co-authored-by: George Muravei-Alkhavoi <georgezagraid@gmail.com>
@BlivionIaG
perf(rdna): share piecewise-graph input buffers across capture sizes
dc9e536 
The RDNA CUDAGraphWrapper cloned every piece's activation inputs for every
capture size and kept strong refs to the outputs, so graph memory grew by
~14 MB per captured token on Flash-Next. A PIECEWISE ladder with prefill
sizes did not fit (OOM in capture at 2048), and every step above 8 tokens
ran fully eager (~235 ms of host launch per step, any chunk size).

For capture sizes >= VLLM_RDNA_CG_SHARED_MIN_TOKENS (default 16; decode
sizes keep the per-size clones), the wrapper now:
- aliases inputs that an earlier piece of the same forward produced
  (their graph-pool address repeats at replay, as upstream relies on), held
  weakly so the static tree doesn't pin pool blocks;
- copies every other token-parallel input (eager splitting-op outputs,
  runner inputs) into one max-size buffer per input slot, shared by all
  sizes (size-sliced views); parameters and non-token-parallel tensors stay
  aliased;
- keeps weak output refs like upstream.

4x V620 TP=4, Flash-Next MTP0, ladder [1..8, 16, 24, 32, 48, 64, 96, 128,
136, 192, 256, 264, 320, 384, 392, 512, 520, 640, 768, 776, 1024, 1032, 1280,
1536, 1792, 2048] (plus VLLM_RDNA_CG_SHARED_MIN_TOKENS=1, GDN projections in
the graph): capture 56 s / 4.2 GiB (default ladder: 5 s / 0.9 GiB).
Host time of a 1024-token step 235 -> ~98 ms. Small mixed steps (a new
short prompt among 6 decoders) max 169-207 -> 74-88 ms. The default
1024-token mixed step is GPU-bound and does not get faster (511 -> 535 ms).
Cells vs the same tip without the ladder (prefill tok/s/req | decode
tok/s/req | TTFT s | ITL med/mean ms):
  1k/512 c=1   1966 | 44.03 | 0.52 | 22.7/22.7  -> 1984 | 44.56 | 0.52 | 22.4/22.4
  1k/512 c=8    519 | 25.35 | 1.97 | 35.8/39.5  ->  495 | 25.39 | 2.07 | 35.5/39.4
  16k/1k c=1   2087 | 43.53 | 7.85 | 23.0/23.0  -> 2075 | 44.04 | 7.90 | 22.7/22.7
  16k/1k c=8    934 | 12.75 | 17.54 | 36.4/78.4 -> 1006 | 12.82 | 16.29 | 35.8/78.0
Greedy probes, PPL (4.81-4.85), prefix, after-load PASS, 0 faults. The
concurrency probe passed on 3 of 4 ladder arms; one failed on a single
margin-1.27 flip at token 28 of prompt 0. The same flip shape showed up
solo-vs-solo without these changes. The recipes keep the
default ladder; the prefill ladder is opt-in via CG_SIZES.
@BlivionIaG
feat(rdna): VLLM_RDNA_GDN_OPAQUE=0 keeps GDN projections in the compi…
0095c9d 
…led graph

On RDNA the whole Qwen GDN layer runs as the opaque qwen_gdn_full_forward
op, so its projections, gated norm and output projection always run eager
(~2.2 ms of host time per layer per prefill step), even inside piecewise
graphs. VLLM_RDNA_GDN_OPAQUE=0 restores the upstream split (only the
attention core is opaque). It needs VLLM_RDNA_CG_SHARED_MIN_TOKENS=1
(the per-size clones of small captures make non-contiguous piece inputs
contiguous, and inductor's stride asserts then fail at size 8). The state
arenas are now allocated for any RDNA cudagraph mode, not only the opaque
path. Default unchanged (opaque).

Flash-Next 2048-token prefill step, host time: 94 -> 80 ms (ladder arm).
Greedy probes, PPL 4.81, prefix, concurrency PASS (ladder + unopaque, two arms).
@BlivionIaG
feat(sched): bound prefill work in steps that carry decodes (RDNA def…
38d2a44 
…ault)

A step that carries prefill tokens freezes every decoding request for the
whole step. Without a cap the scheduler gives that step the full token
budget: on the 27B AWQ dense model (no Mamba-aligned split at TP=4) a mixed
step carries ~2042 prompt tokens and takes 2.0-2.7 s against a 36 ms decode
step; on Flash-Next it carries the 1024-token block (~0.5 s). In the
reverse case (a long prefill running when short requests arrive) the FCFS
running loop let the long prefill take the whole budget every step: the
new requests got 127 of their 128 prompt tokens and then sat behind it,
blocked 7-18 s.

New SchedulerConfig options, off upstream:
- decode_stall_budget_ms (--decode-stall-budget-ms): while requests decode,
  cap the total prefill tokens of a step so that its cost stays within the
  budget. The cost is a line a + b * tokens fitted online to measured
  prefill-carrying step durations (EngineCore feeds them back, both for
  step() and the batch queue; per-bucket medians, so a JIT or TunableOp
  outlier cannot move it). An efficiency floor keeps chunks large enough
  that the per-step cost stays within 25 % of the per-token work, so a
  model with a large per-step cost keeps big chunks (Flash-Next: 1024)
  and one with a small one gets small chunks (27B: ~200-500). Chunks tile
  Mamba state blocks evenly (no short remainder steps).
- decode_time_share (--decode-time-share): optional pure-decode cadence
  paid in wall time (off by default: on Flash-Next it tripled 1k/512 c=8
  TTFT for +60 decode tokens per decoder during a 16k prefill).
With the bound on, running requests with <= 256 tokens left go before long
prefill chunks, and a running prefill leaves up to half the budget for
waiting requests.

The ROCm platform enables it on RDNA (budget 250 ms, share 0;
VLLM_RDNA_DECODE_STALL_MS / VLLM_RDNA_DECODE_SHARE override, an explicit
--decode-stall-budget-ms wins). CUDA/CDNA/other platforms: unchanged.

4x V620 TP=4, HIP 6-9, stall probe (6 decoders, 3 repeats, medians):
27B AWQ   16k inject: max gap 2672 -> 394 ms, gap p50 1978 -> 266 ms,
          tokens/decoder 11 -> 81, x solo 1.03 -> 1.25;
          4k inject: max gap 1972 -> 262 ms, tokens 32 -> 22/dec (136
          total); reverse: decoders blocked 17.8 -> 4.8 s.
Flash-Next (chunks stay 1024): inject unchanged (509 -> 519 ms, 19 ->
          18 tokens); reverse: decoders blocked 7.05 -> 3.18 s, 0 -> 58
          decode tokens during the prefill.
27B cells (prefill tok/s/req | decode tok/s/req | TTFT s | ITL med/mean ms),
before (port notes, same GPUs) -> after:
  1k/512 c=1   1137 | 43.84 | 0.90 | 22.8      -> 1127 | 43.83 | 0.91 | 22.8/22.8
  1k/512 c=8    264 | 23.28 | 3.88 | 36.5      ->  382 | 21.78 | 2.68 | 36.6/45.9
  16k/1k c=1    905 | 38.79 | 18.11 | 25.8     ->  908 | 38.84 | 18.04 | 25.8/25.7
  16k/1k c=8    406 | 6.49 | 40.40 | 52.9      ->  355 | 6.03 | 46.09 | 53.5/165.8
Flash-Next cells, same tree with the bound off (VLLM_RDNA_DECODE_STALL_MS=0) -> on:
  1k/512 c=1   1937 | 26.75 | 0.53 | 22.7/37.4 -> 1956 | 44.00 | 0.52 | 22.7/22.7
  1k/512 c=8    506 | 25.17 | 2.02 | 36.0/39.7 ->  757 | 24.49 | 1.35 | 35.9/40.8
  16k/1k c=1   2042 | 43.49 | 8.02 | 23.0/23.0 -> 2071 | 43.50 | 7.91 | 23.0/23.0
  16k/1k c=8   1042 | 13.02 | 15.72 | 36.3/76.8 -> 1113 | 12.90 | 14.73 | 36.5/77.5
(the off-arm 1k c=1 decode row is a cold-start outlier: ITL med 22.7 both).
Greedy probes, PPL, prefix, after-load and concurrency probes PASS, 0 faults.
@BlivionIaG
feat(rdna): decode-stall cap off by default, enabled in the dense 27B…
5d9f76f 
… recipes

The cap from the previous commit helps dense hybrids (27B AWQ: max
decode gap 2.67 s -> 0.33 s, decoders get 8x more tokens during a 16k
prefill, reverse-case wait 17.9 -> 4.8 s, 1k c=8 TTFT 4.0 -> 3.0 s; cost
16k/1k c=8 TTFT +14 %) and is neutral-to-negative for the MoE models
(Flash-Next inject unchanged, 16k c=8 slower with the ladder). Default
VLLM_RDNA_DECODE_STALL_MS to 0 and set 250 in 27b-awq, 27b-exl3 and full
(27B TP=2); pass the env through serve_rdna.sh.
@BlivionIaG
perf(rdna): fused HC prefill mix kernels for Flash-Next (opt-in)
9cc6d6c 
The prefill (M > 8) hyper-connection mix ran rocBLAS down, a Triton silu,
rocBLAS up writing the [M, 4H] gate, and a Triton gate mix reading it back.
csrc/rocm/hc_prefill_rdna2.cu adds:

- rdna_hc_up_gate_mix_prefill: silu + up GEMM + sigmoid + gated mean in one
  kernel. Each block owns BM tokens x BH hidden columns and computes the 4*BH
  gate columns feeding them, so the gate lives only in registers; silu(x/4)
  is applied while staging dai into LDS (no silu output, no lora copy).
  v_dot2_f32_f16, wave32, LDS k-pair-major tiles, register prefetch;
  22-25 TFLOP/s for the up GEMM it absorbs (rocBLAS tuned: 17-23 for the
  GEMM alone).
- rdna_hc_mix_prefill: split-K HIP down GEMM (N=336, K=10240, fp32 atomic
  accumulator; 1.1-1.2x rocBLAS at M=512..2048, ~1.9x at M<=128) feeding
  the fused kernel, which also writes the fp16 dai for the injection logits.

Both run on the current stream with fresh outputs (graph-capture safe).
Dispatched from rdna_hc_mix with VLLM_RDNA_HC_PREFILL_FUSED=1 (rocBLAS down
+ fused) or 2 (both HIP); default 0 keeps the torch path. Decode (M <= 8)
is unchanged. Ops are registered for gfx1030 builds only.

One HC mix, single V620 (HIP 6), TunableOp rows, graph replay, ms:

     M   torch  mode1  mode2  x mode2
    64   0.140  0.122  0.080  1.74
   128   0.181  0.156  0.156  1.16
   256   0.238  0.203  0.195  1.22
   512   0.426  0.384  0.354  1.20
  1024   0.766  0.710  0.673  1.14
  1030   0.881  0.802  0.761  1.16
  2048   1.468  1.388  1.311  1.12

Max |err| vs torch: mode 1 4.9e-4, mode 2 9.8e-4 (fp16 ulp).
tests/models/qwen4_exp/test_hc_prefill_rdna2.py: 33 passed (M 9..2048 incl.
non-multiples of 64, decode bit-identical, non-contiguous fallback, CUDA
graph replay, shape checks).
@BlivionIaG
perf(rdna): fused HC prefill mix on by default on gfx1030
a553d55 
VLLM_RDNA_HC_PREFILL_FUSED now defaults to 2 (split-K down GEMM + fused
silu/up/sigmoid/mix kernel); 0 restores the torch path. It applies only when
the gfx1030 ops are present and on gfx10x (checked once), otherwise torch.
Decode (M <= 8) is unchanged. serve_rdna.sh prints the knob in its plan.

Flash-Next MTP0, 4x V620 TP=4 (HIP 6-9), F&P, same tree, one boot per arm
(prefill tok/s/req, decode tok/s/req, TTFT s, ITL median / mean ms):

  cell          torch (0)                      fused (2)
  1k/512 c=1    1913 / 15.40* / 0.54 / 22.7 / 64.9   2041 / 43.93 / 0.50 / 22.8 / 22.8
  1k/512 c=8     514 / 25.30 / 1.99 / 35.9 / 39.5     534 / 25.27 / 1.92 / 36.1 / 39.6
  16k/1k c=1    2051 / 43.56 / 7.99 / 23.0 / 23.0    2151 / 43.45 / 7.62 / 23.0 / 23.0
  16k/1k c=8    1051 / 13.04 / 15.59 / 36.5 / 76.7   1089 / 13.29 / 15.05 / 36.6 / 75.2
  (* torch 1k c=1 decode hit an 86 s rdna_ar late-peer stall; ITL median
     equal.) Aggregate output tok/s: 15.2*/184.4/32.5/86.8 -> 42.2/184.8/
  32.9/88.7.

Stall probe (6 decoders, 3 repeats): solo 16k TTFT 7.36 -> 7.15 s, 4k
1.82 -> 1.76 s; inject 16k max gap 504 -> 505 ms, mixed step mean 469 ->
471 ms (the HC saving is ~2 % of a mixed step; the stall floor is elsewhere).
Gates: probe, PPL (4.79 -> 4.81), prefix, concurrency (0 fail; 12/5 ->
9/3 near-tie/soft), probe after load: PASS both; 0 faults.

The 0.28 HIP glue path (VLLM_RDNA_HC_PREFILL_HIP) stays opt-in; its
kernels ignore threadIdx (every thread of a block redoes the whole row), so
silu + gate mix are ~8x slower than Triton (1024 tokens: 5.9 vs 0.76 ms per
HC mix). Docstring updated.
@BlivionIaG
docs(rdna): fused HC prefill mix notes (design, microbench, serve, ol…
42dc168 
…d HIP path)
@BlivionIaG
style(rdna): ruff format the fused HC prefill files
946668b
@BlivionIaG
fix(rdna): decode-stall cap fits mixed steps only; one bucket may shr…
f8af9e1 
…ink the cap

Two reasons the cap never shrank Flash-Next's mixed-step chunks:

1. The line t = a + b * tokens was fitted on every prefill-carrying step.
   Prefill-only steps (deep-context chunks, prompt-logprob steps from the PPL
   probe, prefix-cache tails) are slow for their token count and pulled the
   fit to a = 330-795 ms, b = 0.05-0.28 ms/tok, above the 250/280 ms budget.
   The efficiency floor then pinned the cap at 2048, and mixed steps carried
   whole 1024-token mamba blocks. Only steps with decoders feed the fit now
   (VLLM_DECODE_STALL_FIT_ALL=1 restores the old fit).
2. With a single bucket the cap was max(initial_cap, budget/t*x), so at the
   512 initial cap with ~325 ms steps it could never go lower and never see
   a second bucket. Now it is budget/t*x. Not GPU-tested yet.

Flash-Next MTP0, 4x V620 TP=4 (HIP 6-9), fused HC, CG_SIZES ladder up to
1032, VLLM_RDNA_DECODE_STALL_MS=250, VLLM_DECODE_STALL_MAX_OVERHEAD=1.0,
6 decoders, 3 repeats (fix 1 only; median):

                    before (capB)        after (capC)
  inject 16k        516 ms max gap       341 ms
                    115 dec tok (19/dec) 215 (35/dec)
                    TTFT x solo 1.16     1.52
  inject 4k         504 ms / 37 tok      337 ms / 61 tok
  reverse 16k       3021 ms              3074 ms
  mixed-step chunk  1024 (cap 2048)      512 (cap held at initial 512)
  mean mixed step   477-482 ms           326-328 ms

Gates PASS (probe, PPL, prefix, concurrency, probe after load), 0 faults.
Cells (prefill tok/s/req, decode tok/s/req, TTFT, ITL med/mean):
1k c=1 2011/43.23/0.51/23.1/23.1; 1k c=8 543/23.46/1.89/36.7/42.6;
16k c=1 2114/42.76/7.75/23.4/23.4; 16k c=8 855/11.30/19.16/37.5/88.5
(16k c=8 aggregate output 85.1 -> 74.1 tok/s: smaller chunks cost prefill).
The cap stays off for Flash-Next by default (VLLM_RDNA_DECODE_STALL_MS=0).
@BlivionIaG
tools(rdna): warm every capture size before the stall probe and bench…
7662633 
… cells

The first step at a new token count JIT-compiles / tunes (tens of seconds,
seen as rdna_ar "late peer" warnings and 1k-cell outliers in the ladder
runs: capB 1k c=8 ITL mean 70 vs median 37 ms, decode 14.3 vs 24.5
tok/s/req). serve_validate.sh now runs ladder_warmup.py after the
correctness probes: one exact-length prefill per CG_SIZES entry (fresh random
token ids, so no prefix-cache hit) and decode batches 1..SEQS. WARMUP=0
skips it. Not run on a GPU yet.
@BlivionIaG
fix(rdna): allocate the GDN opaque-forward output in the graph
ab72778 
On RDNA the Qwen GDN layer runs as the opaque qwen_gdn_full_forward op. Its
output was a persistent buffer sized to max_cudagraph_capture_size, sliced
per call and allocated inside the eager-break method. Once the capture
ladder went above 8 tokens, the next compiled piece received the 512-row
buffer while compiled for 16 tokens (warmup died in
assert_size_stride(... 512==16) on 27B AWQ, CG_SIZES up to 512). The buffer
only worked because the RDNA replay-copy CUDAGraphWrapper re-staged piece
inputs.

The output is now a plain empty_like, allocated before the eager break.
Piecewise capture records it in the graph pool, and breakable replay writes
the address the next segment reads, as upstream's output-buffer ops do.

27B AWQ, 4x V620 TP=4, F&P, wrapper removed, CG_SIZES up to 512: boots,
graph memory 1.44 GiB. Probes, PPL 4.94, prefix, concurrency 47/48 identical
PASS. Cells (prefill tok/s/req | decode tok/s/req | TTFT s | ITL med ms):
1k c=1 1111 | 36.84 | 0.92 | 22.9; 1k c=8 257 | 22.59 | 3.99 | 37.8;
16k c=1 888 | 38.58 | 18.44 | 25.9; 16k c=8 401 | 6.39 | 40.89 | 54.2.
The 1k c=1 mean comes from one 8.5 s late-peer stall (first-use JIT on rank 3).
@BlivionIaG
bench(rdna): TP all-reduce microbenchmark at serving sizes + RCCL sweep
2603ff0 
benchmark_rdna_allreduce.py times rdna_ar one-shot, PYNCCL, the CudaCommunicator
dispatch and RS+AG at tokens x hidden fp16 sizes (decode 1-24, mixed 64-1024,
prefill 2048 tokens), eager and replayed from a CUDA graph, through vLLM's own
TP group, and checks every result against the exact sum. ar_sweep.sh runs it
once per environment setting, stopping at the first new PCIe AER or BMC PCI
SERR/PERR record.

4x V620 HIP 2-5 (us/call, eager): RCCL is link-bound from 4 MB (8 MB 501 us,
20 MB 1228 us, 25.6 GB/s bus bw). LL/LL128/Tree/MIN_NCHANNELS do not help above
100 KiB (LL 3-5x slower, Tree 1.4-1.7x; MIN_NCHANNELS is ignored below 8 GPUs).
The one-shot beats RCCL up to ~160 KiB: 80 KiB 47 vs 61, 96 KiB 55 vs 63;
graph 45 vs 65 and 49 vs 69. RCCL wins from 256 KiB. Tables in
docs/rdna2/tp4-allreduce-sizes-2026-10-10.md.
@BlivionIaG
perf(rdna): remove the CUDA-graph replay-copy wrapper, use upstream r…
3491035 
…eplay

The 0.28-era workarounds for HIP graph corruption are gone. The causes were
fixed separately: current-stream launches, fresh op outputs (8b06c28,
00fcc97), mamba slots (5bf02f7), PLE padded replay (cd8056f), and the
GDN output buffer (previous commit).

- CUDAGraphWrapper, make_copy_and_call, caching: upstream again. Gone: per-size
  activation clones, replay copies, strong output refs, NaN scrub,
  shared-buffer mode, staged contiguous inputs.
- Piecewise graphs no longer split at TP collectives. The all-reduce
  (rdna_ar one-shot / RCCL) is captured inside the pieces, as upstream does.
- Breakable eager breaks use upstream weak refs.
- C++: drop rdna2_set_graph_capturing, rdna2_freeze_capture_persist,
  rdna2_immortal_zeros and rdna2_keep_if_capturing. The keepalive header
  becomes rdna2_scratch.cuh: a grow-only split-K scratch for FA-RDNA2 that
  retires grown buffers instead of freeing them.
- Delete the unused vllm/utils/rocm_graph_keepalive.py and the wrapper's
  tests (tests/compile/test_cudagraph_replay_inputs.py).

4x V620 TP=4, F&P. Before = wrapper, after = removed. Cells:
prefill tok/s/req | decode tok/s/req | TTFT s | ITL med ms; graph mem/GPU.
Flash-Next MTP0 (wrapper switched off, same tree):
  1k c=1  1925 | 43.46 | 0.53 | 23.0 -> 1925 | 43.69 | 0.53 | 22.9
  1k c=8   508 | 24.97 | 2.02 | 36.4 ->  502 | 24.90 | 2.04 | 36.4
  16k c=1 2040 | 42.94 | 8.03 | 23.3 -> 2014 | 43.17 | 8.14 | 23.2
  16k c=8 1037 | 12.85 | 15.79 | 37.1 -> 1025 | 12.84 | 15.99 | 36.7
  graphs 0.86 -> 0.13 GiB
27B AWQ:
  1k c=1  1108 | 43.34 | 0.92 | 23.1 -> 1109 | 43.25 | 0.92 | 23.1
  1k c=8   252 | 22.85 | 4.07 | 37.3 ->  258 | 22.87 | 3.96 | 37.2
  16k c=1  888 | 38.23 | 18.46 | 26.2 -> 898 | 38.19 | 18.25 | 26.2
  16k c=8  401 | 6.40 | 40.90 | 53.9 ->  402 | 6.41 | 40.73 | 53.9
  graphs 2.8 -> 0.44 GiB
Ladder up to 512: 27B graphs 13 GiB / c=1 ITL 26.2 ms with the wrapper
(agent S) -> 1.44 GiB / 22.9 ms. Flash-Next 1.10 GiB, cells at parity.

Probes, PPL, prefix and after-load probes PASS, 0 faults, 0 dmesg. One
concurrency flip on the Flash-Next 512 ladder (margin 2.94) did not
reproduce on a fresh boot of the same config. Still pending: the
wrapper-on ladder A/B and more repeats.
@BlivionIaG
docs(rdna): decode-stall cap settings and trade-off; opt-in for Flash…
9239abc 
…-Next

Document the cap's knobs, the mixed-step-only fit, the dense 27B
measurements (on in those recipes) and the Flash-Next arms (opt-in:
max decode gap 505 -> 341 ms for -13 % 16k c=8 output). Pass the
efficiency floor and fit-all switches through serve_rdna.sh and leave a
commented opt-in in the Flash-Next recipes.
@BlivionIaG
style(rdna): ruff format the all-reduce microbenchmark
3b43339
@BlivionIaG
fix(rdna): decode-stall cap fit mode per model: all steps on dense, m…
e09a8c9 
…ixed steps on MoE

The mixed-steps-only fit (f8af9e1) fixed Flash-Next but regressed the
dense 27B AWQ default: the fit drifted up on deep-context mixed chunks and
the 16k injection max decode gap went 394 -> 776 ms. Pick the mode from
model_config.is_moe; VLLM_DECODE_STALL_FIT_ALL=1/0 still forces either.
@BlivionIaG
fix(rdna): keep strong refs in breakable-graph eager breaks on RDNA
dc30658 
The replay-copy wrapper removal (3491035) switched RDNA eager breaks to
upstream weak refs. On RDNA the inductor temps backing GDN/FA inputs are not
always graph-pool owned, so weak refs dangle into capture-time dummy
activations: 27B AWQ TP=2 (full.env, breakable graphs) decoded ' is is is'
after a correct prefill (greedy probe 0/3, PPL fine). Breakable off, or the
pre-removal tree, both pass 3/3. Restore the RDNA strong-ref branch.
@BlivionIaG
tools(rdna): agentic Flash-Next recipe on the flashnext-mtp0 performa…
27861d3 
…nce block

flashnext.env still carried VLLM_RDNA_FUSED_HC=0 and the 0.28 RCCL/alloc
settings, so the agentic launcher ran without the fused hyper-connection
(decode and prefill), tuned MoE configs or the startup-plan/SDMA settings.
Use flashnext-mtp0's env block; keep only the agentic serving deltas.
@BlivionIaG
tools(rdna): decode-stall cap off in full.env (27B TP=2)
2d6e1cd 
At TP=2 each step is ~2x slower, so the 250 ms budget forced tiny 16k
prefill chunks: 16k/1k c=8 aggregate output 22.7 -> 19.6 tok/s, TTFT
135 -> 200 s, ITL median 80 -> 211 ms. The previous plan's no-cap numbers
are the validated path; the TP=4 dense recipes keep the cap.
@BlivionIaG
tools(rdna): dense recipe cap comments carry the 2026-10-10 measurements
e9d22de
@BlivionIaG
tools(rdna): unit.sh covers the fused HC prefill, stall cap, stall pr…
ed4a1c5 
…obe and rdna_ar GPU tests

Drop test_cudagraph_replay_inputs.py (deleted with the replay-copy wrapper).
@BlivionIaG
Revert "tools(rdna): agentic Flash-Next recipe on the flashnext-mtp0 …
0272b15 
…performance block"

This reverts commit 27861d3. The aligned recipe failed its validation
(2026-10-11 01:15): during the prefix probe several workers aborted with
'LLVM ERROR: out of memory' in a runtime inductor compile of the shared-expert
path, killing the engine. Back to the recipe that passed the final pass
(fused hyper-connection off) until the aligned one is debugged.
@BlivionIaG
test(rdna): GDN arena test follows the decode-batch sizing (ad356d1)
076fd4d 
The V2 runner never captures a decode graph with more requests than
max_num_seqs, so the arenas are sized by max_num_seqs * (num_spec + 1)
whatever the capture ladder. Check that, including a 2048 ladder, instead of
the old capture-size expectation. 53/53 on gfx1030.
@BlivionIaG
tools(rdna): agentic Flash-Next recipe on the flashnext-mtp0 performa…
20edef5 
…nce block (again)

Re-applies 27861d3. Its first validation hit a one-off 'LLVM ERROR: out of
memory' in a runtime compile; the bisect (2026-10-11) passed every gate three
times: the aligned recipe unchanged, the old recipe with only the fused
hyper-connection, and the aligned recipe with the compile cache off.
@BlivionIaG
BlivionIaG changed the base branch from rdna_extra/v0.28.0 to rdna_extras October 11, 2026 01:59
@BlivionIaG
BlivionIaG merged commit 7813953 into rdna_extras Oct 11, 2026
2 checks passed
@BlivionIaG BlivionIaG mentioned this pull request Oct 11, 2026
Merged
This was referenced Oct 11, 2026
Draft
Draft
@BlivionIaG BlivionIaG mentioned this pull request Oct 16, 2026
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

5 participants

@BlivionIaG @GeorgeMurAlkh @leapdragon @Karl0007 @wsantos