Skip to content

[review-only] cherry-pick leapdragon rdna_ar fabric knobs onto rdna_extras - #1

Merged
BlivionIaG merged 5 commits into
rdna_extrasfrom
cursor/leapdragon-cherry-d2a4
Sep 8, 2026
Merged

[review-only] cherry-pick leapdragon rdna_ar fabric knobs onto rdna_extras#1
BlivionIaG merged 5 commits into
rdna_extrasfrom
cursor/leapdragon-cherry-d2a4

Conversation

@BlivionIaG

@BlivionIaG BlivionIaG commented Sep 3, 2026 •
edited by cursor Bot
Loading

Copy link
Copy Markdown

Purpose

Review only. This is not dest. Do not merge as dest. Occupancy pin stays closed. Compiler pin stays ROCm 7.14.

Side-branch cherry-picks of HIP custom allreduce / rdna_ar fabric knobs from leapdragon/vllm-rdna2-qwen branch rdna2/qwen38-flash-next onto rdna_extras for UNC-27 review.

Authorship is theirs. Every unique commit is Author and Committer Aron Hsiao <leapdragon@gmail.com>. Original Co-Authored-By: Claude Fable 5 trailers from the source commits are kept. This branch does not add BlivionIaG, Cursor, or any other local/agent trailer to their work.

AR default stays off. Dest policy is opt-in (VLLM_FORCE_CUSTOM_ALL_REDUCE and VLLM_RDNA_AR). Wiring uses os.getenv("VLLM_RDNA_AR", "0") == "1". Do not treat this path as dest-on-by-default.

Kept: hipExtMallocWithFlags(..., hipDeviceMallocUncached) + peer push. No Finegrained alloc. No WaitValue-in-capture. Dest HIP (packed fdot2 / FA-RDNA2 / EXL3 3inst) was not replaced. Their gemv_f16_rdna2 / moe_skinny_int4_decode were dropped at conflict resolution.

Source

  • Repo: leapdragon/vllm-rdna2-qwen
  • Branch: rdna2/qwen38-flash-next

Picked commits (Author + Committer: Aron Hsiao <leapdragon@gmail.com>)

Cherry-pick -x (Author preserved; not a squash).

Source SHA Side-branch SHA Subject
fe2b63f37061741c40af678921e2b71c5b46640c af25c53292111bdc93ed9431d8523ca3fd7181c7 T44: push-based one-shot all-reduce (Uncached+IPC)
b84ea91e052b71453670a90843c1cda65d69576f f612cf1fab1abf653ca2e7d02fa4103b74646cdb fabric-friendlier pushes (stagger / VLLM_RDNA_AR_BLOCKS / VLLM_RDNA_AR_PACE)
3a4cf504705df01c19ad23a3908867484f98229b 351ee49684da1abd71aae8cdf958bf22fcb75204 fix log-line syntax
5e93db1f0a79cb3971899cff515cef4c82ece91e 0f075689e5da68cfd6156831587d8d0d09061818 boot self-test (PLE docs not taken)
2b703492b32388c669478082a4ce694bf740a7c0 ee6e48ea111eb9aeed0a6c3d689ecac94e25e15e boot self-test warmup (min of 3 timed repeats)

Conflict resolutions (AR-only wiring; kernels not rewritten):

  • T44: dest ops.h / _custom_ops.py / cuda_communicator.py — add rdna_ar_* only; keep dest HIP; all_reduce() keeps dest order (symm-mem / qr / FI first); rdna_ar is a later opt-in path.
  • b84ea91: launch args keep T44 g.dflags (host-coherent flags) and add g.pace. Did not take 3cfe000 flag-in-VRAM rewrite (that commit is PLE-mixed).
  • 5e93db1: kept rdna_all_reduce.py self-test; dropped docs/rdna2/PLE-*.md dest does not have.

Not picked (denylist / mixed)

  • 3cfe00035559a1bff0b05b6ac82cac5ad08fd5bc — uncached VRAM flags plus PLE worker/tests
  • 5d54e58f061e83e8150d02c9c67c4ba89e0a7fad — T45/T46 mixed int8 / MoE / gemv_f16
  • 29305a9f858386a7020ec592af53e1c8258dd51f — PLE + tok/s docs
  • 11c04dffabcabe33d9428bd0eece5313a0f7cac1 — docs with their measured µs/op (not dest numbers)
  • a600a1f7846deb024a9f41b5f7e65f7b81d70272 — PLE docs only
  • containers, ROCR idle-CPU, PLE/MTP/n-gram, recipe 0006, wvSplitK_int4, scalar nibble MoE, dest HIP replacements

Why this is not duplicating an existing PR

No other open dest PR was found for this rdna_ar allowlist (rdna_ar / leapdragon search on opengfx1030/vllm-rdna).

Test Plan

  • Cherry-pick -x on side branch only; rdna_extras tip must stay f93619500c3e55daee5f60ae7f6d4279cdeb3f0d.
  • Verify Author and Committer of all unique commits are Aron Hsiao; no BlivionIaG / Cursor trailers.
  • Verify VLLM_RDNA_AR defaults to "0"; dest FA-RDNA2 / EXL3 / GDN decls remain; hipDeviceMallocUncached + push remain.
  • No GPU e2e / model eval in this environment (review-only). Hardware check for UNC-27 is follow-up.

Test Result

  • Five unique commits on cursor/leapdragon-cherry-d2a4, Author and Committer Aron Hsiao.
  • origin/rdna_extras still f93619500c3e55daee5f60ae7f6d4279cdeb3f0d.
  • No model evaluation (opt-in AR wiring + new rdna_ar sources; dest HIP kernels unchanged).
Essential Elements of an Effective PR Description Checklist
  • The purpose of the PR (review-only cherry-pick; not dest).
  • The test plan (cherry-pick / Author / default-off / dest-tip checks).
  • The test results (5 Aron Hsiao commits; dest tip unchanged; no GPU eval).
  • Documentation update (not dest; no tok/s copied in as dest numbers).
Open in Web Open in Cursor 

leapdragon and others added 5 commits August 29, 2026 14:48
@leapdragon
T44: push-based one-shot all-reduce for small TP messages on gfx1030,…
af25c53 
… W=2..8

rdna_ar_oneshot (csrc/rocm/rdna_allreduce.cuh): each rank pushes its contribution into
every peer's uncached staging buffer, signals via host-coherent flags, reduces in fixed
rank order (bit-identical across ranks); device-derived sequence numbers so CUDA-graph
replay is safe; bounded spins. Per-process-group instances (handles). Hooked into
CudaCommunicator.all_reduce ahead of every other path; VLLM_RDNA_AR=0 disables.
4x V620 cross-PCIe: 33 us/op for the 20 KB decode message (RCCL in-server: 156 us).

(cherry picked from commit fe2b63f)
@leapdragon
rdna_ar: fabric-friendlier pushes — rank-staggered peer order, VLLM_R…
f612cf1 
…DNA_AR_BLOCKS cap, VLLM_RDNA_AR_PACE inter-store idle

At any instant each destination GPU now receives from one source instead of all W-1 (pure
reorder, no cost); the block cap and pace bound the concurrent PCIe push streams and their
rate into the receiving root complex. Defaults unchanged (auto blocks, pace 0).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit b84ea91)
@leapdragon @claude
rdna_ar: fix log-line syntax
351ee49 
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 3a4cf50)
@leapdragon
rdna_ar: boot self-test (verified collectives, timeout + latency boun…
0f07568 
…d, group-wide verdict)

On boards with weak GPU P2P (ACS redirect, chipset-routed or cross-socket slots) init
could succeed while every collective then spun to its ~2 s cap and aborted WITHOUT
writing the output -- silent corruption reported by whatever waits next (usually the
PLE handshake). Three known-pattern all-reduces now run at init; wrong results, a
sticky timeout, or >50 ms latency disable the fast path group-wide with a clear log
line ('rdna_ar: disabled -- boot self-test failed'), falling back to RCCL. Docs:
diagnostic branch G in PLE-OFFLOAD-SETUP.md / PLE-DIAGNOSTIC-TREE.md.

Verified on 4x V620: self-test passes, fast path active. Noted: one 600-token
determinism check after this change showed a divergence (not investigated per
operator's call; the active instance reported handle 1 where earlier boots showed
handle 0).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 5e93db1)
@leapdragon @claude
rdna_ar: boot self-test warms up each size and judges the minimum of …
ee6e48e 
…3 timed repeats

The single timed collective per size included rank skew at boot: a healthy machine failed the
50 ms bound on trial 1 (59 ms) right after a fresh reboot and silently fell back to RCCL (-25 %
decode). Warm-up absorbs the first-launch code-object load and IPC first-touch; the minimum is
what a genuinely slow P2P path cannot hide.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 2b70349)
@cursor
cursor Bot force-pushed the cursor/leapdragon-cherry-d2a4 branch from 4ec3ffe to ee6e48e Compare September 3, 2026 13:44
@cursor cursor Bot mentioned this pull request Sep 8, 2026
Merged
@BlivionIaG
BlivionIaG marked this pull request as ready for review September 8, 2026 21:57
@BlivionIaG
BlivionIaG merged commit a406064 into rdna_extras Sep 8, 2026
@BlivionIaG BlivionIaG mentioned this pull request Sep 9, 2026
Merged
BlivionIaG added a commit that referenced this pull request Sep 18, 2026
@BlivionIaG @leapdragon @claude
[review-only] cherry-pick leapdragon rdna_ar fabric knobs onto rdna_e…
383b418 
…xtras (#1)

* T44: push-based one-shot all-reduce for small TP messages on gfx1030, W=2..8

rdna_ar_oneshot (csrc/rocm/rdna_allreduce.cuh): each rank pushes its contribution into
every peer's uncached staging buffer, signals via host-coherent flags, reduces in fixed
rank order (bit-identical across ranks); device-derived sequence numbers so CUDA-graph
replay is safe; bounded spins. Per-process-group instances (handles). Hooked into
CudaCommunicator.all_reduce ahead of every other path; VLLM_RDNA_AR=0 disables.
4x V620 cross-PCIe: 33 us/op for the 20 KB decode message (RCCL in-server: 156 us).

(cherry picked from commit fe2b63f)

* rdna_ar: fabric-friendlier pushes — rank-staggered peer order, VLLM_RDNA_AR_BLOCKS cap, VLLM_RDNA_AR_PACE inter-store idle

At any instant each destination GPU now receives from one source instead of all W-1 (pure
reorder, no cost); the block cap and pace bound the concurrent PCIe push streams and their
rate into the receiving root complex. Defaults unchanged (auto blocks, pace 0).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit b84ea91)

* rdna_ar: fix log-line syntax

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 3a4cf50)

* rdna_ar: boot self-test (verified collectives, timeout + latency bound, group-wide verdict)

On boards with weak GPU P2P (ACS redirect, chipset-routed or cross-socket slots) init
could succeed while every collective then spun to its ~2 s cap and aborted WITHOUT
writing the output -- silent corruption reported by whatever waits next (usually the
PLE handshake). Three known-pattern all-reduces now run at init; wrong results, a
sticky timeout, or >50 ms latency disable the fast path group-wide with a clear log
line ('rdna_ar: disabled -- boot self-test failed'), falling back to RCCL. Docs:
diagnostic branch G in PLE-OFFLOAD-SETUP.md / PLE-DIAGNOSTIC-TREE.md.

Verified on 4x V620: self-test passes, fast path active. Noted: one 600-token
determinism check after this change showed a divergence (not investigated per
operator's call; the active instance reported handle 1 where earlier boots showed
handle 0).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 5e93db1)

* rdna_ar: boot self-test warms up each size and judges the minimum of 3 timed repeats

The single timed collective per size included rank skew at boot: a healthy machine failed the
50 ms bound on trial 1 (59 ms) right after a fresh reboot and silently fell back to RCCL (-25 %
decode). Warm-up absorbs the first-launch code-object load and IPC first-touch; the minimum is
what a genuinely slow P2P path cannot hide.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 2b70349)

---------

Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
BlivionIaG added a commit that referenced this pull request Sep 21, 2026
@BlivionIaG @leapdragon @claude
[review-only] cherry-pick leapdragon rdna_ar fabric knobs onto rdna_e…
e29e0cb 
…xtras (#1)

* T44: push-based one-shot all-reduce for small TP messages on gfx1030, W=2..8

rdna_ar_oneshot (csrc/rocm/rdna_allreduce.cuh): each rank pushes its contribution into
every peer's uncached staging buffer, signals via host-coherent flags, reduces in fixed
rank order (bit-identical across ranks); device-derived sequence numbers so CUDA-graph
replay is safe; bounded spins. Per-process-group instances (handles). Hooked into
CudaCommunicator.all_reduce ahead of every other path; VLLM_RDNA_AR=0 disables.
4x V620 cross-PCIe: 33 us/op for the 20 KB decode message (RCCL in-server: 156 us).

(cherry picked from commit fe2b63f)

* rdna_ar: fabric-friendlier pushes — rank-staggered peer order, VLLM_RDNA_AR_BLOCKS cap, VLLM_RDNA_AR_PACE inter-store idle

At any instant each destination GPU now receives from one source instead of all W-1 (pure
reorder, no cost); the block cap and pace bound the concurrent PCIe push streams and their
rate into the receiving root complex. Defaults unchanged (auto blocks, pace 0).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit b84ea91)

* rdna_ar: fix log-line syntax

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 3a4cf50)

* rdna_ar: boot self-test (verified collectives, timeout + latency bound, group-wide verdict)

On boards with weak GPU P2P (ACS redirect, chipset-routed or cross-socket slots) init
could succeed while every collective then spun to its ~2 s cap and aborted WITHOUT
writing the output -- silent corruption reported by whatever waits next (usually the
PLE handshake). Three known-pattern all-reduces now run at init; wrong results, a
sticky timeout, or >50 ms latency disable the fast path group-wide with a clear log
line ('rdna_ar: disabled -- boot self-test failed'), falling back to RCCL. Docs:
diagnostic branch G in PLE-OFFLOAD-SETUP.md / PLE-DIAGNOSTIC-TREE.md.

Verified on 4x V620: self-test passes, fast path active. Noted: one 600-token
determinism check after this change showed a divergence (not investigated per
operator's call; the active instance reported handle 1 where earlier boots showed
handle 0).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 5e93db1)

* rdna_ar: boot self-test warms up each size and judges the minimum of 3 timed repeats

The single timed collective per size included rank skew at boot: a healthy machine failed the
50 ms bound on trial 1 (59 ms) right after a fresh reboot and silently fell back to RCCL (-25 %
decode). Warm-up absorbs the first-launch code-object load and IPC first-touch; the minimum is
what a genuinely slow P2P path cannot hide.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
(cherry picked from commit 2b70349)

---------

Co-authored-by: Aron Hsiao <leapdragon@gmail.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
jmacpolin pushed a commit to jmacpolin/vllm-rdna that referenced this pull request Oct 13, 2026
@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 (opengfx1030#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>
jmacpolin pushed a commit to jmacpolin/vllm-rdna that referenced this pull request Oct 13, 2026
@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 (opengfx1030#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 opengfx1030#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) (opengfx1030#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>
jmacpolin pushed a commit to jmacpolin/vllm-rdna that referenced this pull request Oct 13, 2026
@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 (opengfx1030#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 opengfx1030#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) (opengfx1030#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 opengfx1030#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>
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.

2 participants

@BlivionIaG @leapdragon