From b8b3b65b2c3616bca322501e73a67b7f19f948f1 Mon Sep 17 00:00:00 2001 From: Cursor Agent Date: Thu, 1 Oct 2026 16:03:12 +0000 Subject: [PATCH] Lock dest extras tracking to rdna_extras @ bc5fbee5 Observe dest serve_rdna recipes, EXL3 quantized MTP draft, TunableOp profiles, and causal_conv HIP state_len gate. Stay extras. Still no torch.ops.hippihx. Co-authored-by: BlivionIaG --- CONTRIBUTING.md | 5 +- docs/ARCHITECTURE.md | 8 ++- docs/BACKPORT.md | 97 +++++++++++++++++-------- docs/CONSUME.md | 8 ++- docs/EXTRAS.md | 102 ++++++++++++++++++--------- tests/test_no_b12x_cuda.py | 67 +++++++++++++++++- tiles/attention/fa_fdot2/README.md | 5 +- tiles/gemm/exl3_3inst/README.md | 8 ++- tiles/sequence/causal_conv/README.md | 2 +- 9 files changed, 227 insertions(+), 75 deletions(-) diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 5d07bb8..05f6120 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -48,8 +48,9 @@ those commits. **Foreign HIP** (leapdragon, a17t, anyone who is not this lab) keeps **their** Author **and** Committer. Same rule as extras [PR #1](https://github.com/opengfx1030/vllm-rdna/pull/1) (leapdragon -`rdna_ar`, now dest extras @ `a4060647` squash / tip `30632b2`; pick unique Aron Hsiao commits, -not the GitHub squash). Do not rewrite those as Blivion or Cursor. +`rdna_ar`; dest squash `a4060647` / tip `bc5fbee5`; pick unique Aron +Hsiao commits, not the GitHub squash). Do not rewrite those as Blivion +or Cursor. | Rule | Meaning | |---|---| diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md index ea9c7e7..b169939 100644 --- a/docs/ARCHITECTURE.md +++ b/docs/ARCHITECTURE.md @@ -41,7 +41,7 @@ Same *shape* as b12x (`.` + `api.py`), HIP objects: | `hippihx/flydsl/` | FlyDSL compiler atoms + kernel contracts | | `hippihx/comm/fabric.py` | PCIe/PLX hop class (PIX/PXB/PHB × 88096/8749) | -Group rename **`attn` → `attention`** (V1 ABI rev **3**; ids unchanged). ISA +Group rename **`attn` → `attention`** (V1 ABI rev **3**; ids unchanged; rev **4** is the consume ABI). ISA class names stay (`fa_fdot2`, not b12x `paged`). Map: | hippihx | b12x analogue | @@ -231,8 +231,10 @@ optional activation `dtype` (revision **2**). Qualnames `attention.*` plan, params, tensor descriptors, the stream and `caps.fabric`. Per-op params, tensor slots and scratch rules live in the catalog, and `python -m hippihx._lib.codegen` writes the header enums and C tables. -See [`CONSUME.md`](CONSUME.md). Dest extras tip `30632b2fa323` still -has no `torch.ops.hippihx.*` rewire — see [`BACKPORT.md`](BACKPORT.md). +See [`CONSUME.md`](CONSUME.md). Dest extras tip `bc5fbee5d0a1` still +has no `torch.ops.hippihx.*` rewire. Dest extras **#27** is opt-in +`VLLM_HIPPIHX` ctypes consume (default **off**; plans not ready) — see +[`BACKPORT.md`](BACKPORT.md). Do not edit `opengfx1030/vllm-rdna` from this tree. ## Non-goals (room lock) diff --git a/docs/BACKPORT.md b/docs/BACKPORT.md index 3767f16..87b6bdf 100644 --- a/docs/BACKPORT.md +++ b/docs/BACKPORT.md @@ -1,10 +1,11 @@ # extras → hippihx backport review Lock: [`opengfx1030/vllm-rdna`](https://github.com/opengfx1030/vllm-rdna) -`rdna_extras` @ `30632b2fa323` (2026-09-27 13:05 UTC). Dest **default +`rdna_extras` @ `bc5fbee5d0a1` (2026-10-01 13:08 UTC). Dest **default branch is `rdna_extras`**, not `main` (`c00091e02670` upstream vLLM — no -dest HIP). **No** `torch.ops.hippihx.*`. Observe only: no `.cu` dump, no -tok/s, no extras maxdiff. +dest HIP). **No** `torch.ops.hippihx.*`. Dest extras **#27** is opt-in +`VLLM_HIPPIHX` ctypes consume (default **off**; plans not ready). Observe +only: no `.cu` dump, no tok/s, no extras maxdiff. hippihx is the zoo. extras is serve wiring. Dest HIP is still ATen-coupled. V1 C ABI (`include/hippihx/v1.h`, rev **4**, @@ -45,21 +46,47 @@ until a body migrates and extras binds it. | `f3dd65fa7063` | extras PR **#15** merge (QSA live-context bound + folded **#12** PLE/MTP/graph-redirect) | Stay extras. Python/Triton/serve. Foreign George + Codex. Do not pick. | | `e131562936f9` | extras PR **#17** merge (resident W4A16 MoE / Mamba #55450 retire) | Stay extras. ATen HIP. `VLLM_RDNA_MOE_RESIDENT*` default **off**. Foreign George + Codex. Do not pick. | | `48c56efbe80b` | pin `AttentionConfig.backend = RDNA_ATTN` when `VLLM_USE_RDNA2_FA=1` | Stay extras. Python platform. Unique Blivion. | -| `68a635ed8d78` | extras **#20** capture + dest-integrated **#21** FULL+PIECEWISE | Stay extras. Keeps FULL decode graphs; piecewise mixed/prefill. Foreign Waldecir + Cursor. Zoo does not own graph mode. | +| `68a635ed8d78` | extras PR **#20** capture + dest-integrated PR **#21** FULL+PIECEWISE | Stay extras. Keeps FULL decode graphs; piecewise mixed/prefill. Foreign Waldecir + Cursor. Zoo does not own graph mode. | | `2a5e89368272` | extras PR **#19** merge (MTP unquantized-weight detect) | Stay extras. Python only. MTP still not dest. Foreign a17t + opencode. Do not pick. | | `700753d9add5` | MTP draft unbreak (amdsmi shutdown swallow + local-argmax) | Stay extras. Python platform + MTP model. MTP still not dest. Unique Blivion. | | `30632b2fa323` | draft decode graph capture for `max_num_reqs` | Stay extras. Python worker (`cudagraph_utils`). MTP still not dest. Unique Blivion. Zoo does not own graph mode. | +| `cd38a1d38cbd` | extras PR **#24** merge (RAM KV offload) | Stay extras. Python serve. Foreign George + Codex. Do not pick. No HIP. | +| `22d6e34625e4` | extras PR **#22** dest-presence (two-shot `rdna_ar`) | Stay extras. ATen HIP. Cursor dest-presence. Do not dump. Do not pick. Dest @ `101a16c8` disabled two-shot by default. | +| `a3f7e5da97f0` | extras PR **#27** merge (opt-in hippihx V1 consume) | Stay extras. Python ctypes. `VLLM_HIPPIHX` default **off**. Plans not ready. Foreign Claude. Do not pick. Still no `torch.ops.hippihx`. | +| `7434efee40cb` | extras PR **#26** merge (HIP MoE skinny `MAX_M=16` opt-in) | Stay extras. Default still **8**. Foreign George + Codex. Do not dump. Do not pick. | +| `101a16c8bef7` | dest-disable two-shot `rdna_ar` by default | Stay extras. Library `MAX_KB` **64**. `VLLM_RDNA_AR` still **0**. Zoo **512**. Do not dump. | +| `330b42abb5b0` | MAXBAT=2048 serve default | Stay extras. Serve scheduler. Do not copy tok/s. | +| `56744dd437da` | launcher safety (full.sh `MAX_KB` **64**, scoped cleanup) | Stay extras. Serve scripts. Do not copy tok/s. | +| `bfd5286d846b` | extras PR **#28** dest-presence (FA tile skip / in-place out / GQA softmax) | Stay extras. ATen HIP. Foreign Claude. Do not dump. Do not pick. | +| `83e6af80d4bf` | dest FA verify-through-decode + `UNIFORM_BATCH` + MTP launcher `ATTN=fa` | Stay extras. Python serve. No HIP. Unique Blivion. Superseded as dest tip by extras **#29** dest-presence. Zoo does not own graph mode. MTP still not dest. Do not copy tok/s. | +| `ae5bedfd61fc` | production launcher PYTHONPATH pin + FA defaults | Stay extras. Serve scripts. Do not copy tok/s. | +| `6ed390937941` | prefill step cadence for non-DP engines | Stay extras. Serve/engine. Co-authored wsantos. Do not pick. Do not copy tok/s. | +| `d94e22099bed` | extras PR **#29** dest-presence (FA split decode / `cu_query_lens`) | Stay extras. ATen HIP. Dest rework of foreign Claude. Do not dump. Do not pick Claude. | +| `e0112c55e48d` | dest split-decode row gate (≤256 query×head) | Stay extras. Python serve. Unique Blivion. Zoo does not own graph mode. Do not copy tok/s. | +| `3a0786eaa70f` | dest W4A8 sdot4 opt-in + K_STEP prefill split repair | Stay extras. ATen HIP. Unique Blivion. Default **off** (`VLLM_RDNA2_W4A8_SDOT4` unset = W4A16). Do not dump. Do not copy tok/s. | +| `3fb9d43afdb5` | extras **#30** dest-presence (W4 compile-dispatch opt-in) | Stay extras. Python only. Unique Blivion dest merge. Default **off** (`VLLM_RDNA2_W4A16_RUNTIME_DISPATCH`). GitHub PR closed-unmerged. Do not dump. Do not copy tok/s. | +| `9cad8792ad89` | dest MoE epilogue CAS default; fp32 scratch opt-in | Stay extras. ATen HIP (`moe_accum_rdna2.cuh`). Unique Blivion. Default **off** (`VLLM_RDNA2_MOE_FP32_ACCUM`). Packed-fp16 CAS is dest default. Do not dump. Do not copy tok/s. | +| `a5059cf0480e` | extras **#32** dest-presence (EXL3 mul1 / K=1..8 + MoE staging bound) | Stay extras. ATen HIP. Unique Blivion dest landing. GitHub PR closed-unmerged. Do not dump. Do not copy tok/s. | +| `30b0bd4e3604` | dest EXL3 v0.30 loader + fused project + MoE method + PLE n-gram | Stay extras. Python + dest HIP `exl3_project_rdna2`. Unique Blivion. extras **#32** GitHub closed-unmerged. Do not dump. Do not copy tok/s. | +| `ac9dfd8fba6d` | dest TunableOp unify + EXL3 27B rows + gemma_rms fake | Stay extras. Serve scripts / rocBLAS tables / inductor fake. Unique Blivion. Do not dump. Do not copy tok/s. | +| `bc5fbee5d0a1` | dest serve_rdna recipes + EXL3 MTP draft + TunableOp profiles | Stay extras. Python/serve. Unique Blivion. Dest now serves EXL3 quantized MTP; no zoo MTP tile. Do not dump. Do not copy tok/s. | Other dest-fixed ISA (keep in tile locks, not a dump): GDN prefill `o` `i_t_local`; causal_conv FIR pre-shift then shift; ConfigH (`K_STEP=64`) reverted. Live dest bugs: [Dest extras defects](#dest-extras-defects). -Open extras PRs **#2** (GLM Later), **#18**, **#22**, **#24**. Draft **#9/#10** -sdot, **#23** PCIe P2P KV — skip. Closed **#3/#16** unmerged, **#21** dest-integrated +Open extras PRs **#2** (GLM Later), **#18**. Draft **#10** sdot8, +**#23** PCIe P2P KV, **#25** rdna_ar retained-output, **#31** W4 +exact-dequant, **#33** FA decode scores, **#34** FA D=128 prefill — +skip. Closed **#3/#9/#16/#30/#32** unmerged (**#9** superseded by dest +W4A8 @ `3a0786ea`; **#30** dest-landed independently @ `3fb9d43a`; +**#32** dest-landed independently @ `a5059cf0`), **#21** dest-integrated unmerged, **#11** dest-reverted wvSplitK (**no zoo tile**), **#12** -superseded by dest **#15**. Merged **#13/#14/#15/#17/#19/#20** dest -*presence*, observe, do not pick. -`rdna_extras_wip_20260910` is not dest. +superseded by dest **#15**. Merged +**#13/#14/#15/#17/#19/#20/#22/#24/#26/#27/#28/#29** dest *presence*; +dest **#30** @ `3fb9d43a` / dest **#32** @ `a5059cf0` dest-presence. +Dest tip `bc5fbee5` is serve recipes + EXL3 MTP (not a new extras PR). +Observe, do not pick. `rdna_extras_wip_20260910` is not dest. ## Action @@ -69,34 +96,44 @@ superseded by dest **#15**. Merged **#13/#14/#15/#17/#19/#20** dest | FA GQA-subgroup prefill + register-resident O | **Later** `attention/fa_fdot2`. Occupancy pin closed. | | causal_conv out-stride + null-block | **Later** `sequence/causal_conv`. | | GDN decode fp16 SSM state | **Later** `attention/gdn_scan`. Recurrence stays 16 fp32 VGPR. | -| leapdragon `rdna_ar` + PIX helpers + dest T44b (`3b59ee16`) | **Later** `comm.pcie`. Still opt-in. Dest extras `MAX_KB` **64**; zoo **512**. Do not pick squash. | +| leapdragon `rdna_ar` + PIX helpers + dest T44b (`3b59ee16`) | **Later** `comm.pcie`. Still opt-in. Dest extras `MAX_KB` **64** (two-shot dest-off @ `101a16c8`). Zoo **512**. Do not pick squash. | | PR **#2** GLM-5.3 KDA/DSA | **Later** `kda_scan` / `dsa_nope` / `qsa_indexer`. Do not name tiles `glm5_*`. | -| GDN arenas, `rdna2_graph_keepalive.cuh`, breakable cudagraphs, Hybrid W4 gfx10, Flash-Next HC/QSA/PLE/M-RoPE HIP, skinny GEMM / `gemv_f16_rdna2`, PLE schema, W4 MoE oracle, `new_zeros`/`zeros_like`, V1 FULL→PIECEWISE, vLLM custom AR, `bench_report.py`, seq cap 6, Flash-Next launcher knobs, HC `_contig()` cache, QSA Triton bounds, mamba spec-decode, T44b wedge, V620 MoE JSON / amdsmi / PLE, ROCm platform init, extras EXL3 docker arch-guard, Qwen4Exp MTP, recovered extras probes, resident MoE, FA RDNA_ATTN pin, `qwen4_exp/**` | **Stay extras.** Product gates default off (S6 revert). No new V1 op until dest locks a class. | -| PR **#3** a17t / explore **#9/#10** sdot / closed **#12** / closed **#16** / PR **#18** / PR **#22** / PR **#23** / PR **#24** | **Skip.** Second W4 family, not dest, serve-only. **#16** unmerged. **#18** GPTQ `BLOCK_KN_SIZE` 256. **#22** two-shot `rdna_ar` (Cursor; library still **0**; zoo **512**; do not dump). **#23** PCIe P2P KV (Cursor; HIP IPC SDMA; zoo `comm.pcie` is AR, not KV disagg). **#24** RAM KV offload (George/Codex; no HIP). | -| PR **#5** / **#8** Flash-Next, extras **#11** wvSplitK, extras **#13** T44b, extras **#14**, extras **#15**, extras **#17**, extras **#19**, extras **#20**, extras **#21** | **Closed.** Dest-landed T44b / #14 / #15 / #17 / #19 / #20 / dest-integrated #21 are observe-only. Dest reverted wvSplitK (`5c4ab989`). Do not re-merge. Do not grow `gemv_f16`. | +| Serve, product HIP, dest-landed extras PRs, MTP, graph mode, `new_zeros`/`zeros_like`, `qwen4_exp/**` | **Stay extras.** Product gates default off (S6 revert). Dest **#27** opt-in ctypes consume; plans not ready. Dest **#28** / dest **#29** split decode / dest row-gate (`e0112c55`) / dest W4A8 (`3a0786ea`) / dest **#30** compile-dispatch (`3fb9d43a`) / dest MoE epilogue (`9cad8792`) / dest **#32** EXL3 mul1 (`a5059cf0`) / dest EXL3 loader (`30b0bd4e`) / dest TunableOp (`ac9dfd8f`) / dest EXL3 MTP + `serve_rdna` (`bc5fbee5`) stay extras; do not dump. No body migrate. See SHA index. | +| PR **#3** a17t / closed **#9** / explore **#10** sdot8 / closed **#12** / closed **#16** / PR **#18** / PR **#23** / PR **#25** / PR **#31** / PR **#33** / PR **#34** | **Skip.** Second W4 family, not dest, serve-only. **#9** closed-unmerged (superseded by dest W4A8). **#16** unmerged. **#18** GPTQ `BLOCK_KN_SIZE` 256. **#23** PCIe P2P KV (HIP IPC SDMA; zoo `comm.pcie` is AR, not KV disagg). **#25** rdna_ar retained-output (George/Codex). **#31** W4 exact-dequant explore. **#33** FA decode scores. **#34** FA D=128 prefill. | +| PR **#5** / **#8** Flash-Next, extras **#11** wvSplitK, extras **#13** T44b, extras **#14**, extras **#15**, extras **#17**, extras **#19**, extras **#20**, extras **#21**, extras **#22**, extras **#24**, extras **#26**, extras **#27**, extras **#28**, extras **#29** | **Closed.** Dest-landed observe-only. Dest reverted wvSplitK (`5c4ab989`). Do not re-merge. Do not grow `gemv_f16`. | Bodies stay in extras because every dest HIP file includes `torch/all.h`. A dump is not a migrate. CONTRIBUTING: do not edit extras from this repo. ## Dest file → tile -Observed at `30632b2fa323`. Numbers are extras observations, not dest +Observed at `bc5fbee5d0a1`. Numbers are extras observations, not dest locks. Fill tile READMEs. | extras path | hippihx tile | Notes | |---|---|---| -| `csrc/rocm/fa_rdna2.cu` | `attention/fa_fdot2` | ~33 KiB decode / ~48 KiB prefill smem. GQA-subgroup default. O accumulator register-resident @ `d1b200b1`. Dest @ `48c56ef` pins `RDNA_ATTN` from API-server `VLLM_USE_RDNA2_FA`. Persist O stay extras. | +| `csrc/rocm/fa_rdna2.cu` | `attention/fa_fdot2` | ~33 KiB decode / ~48 KiB prefill smem. GQA-subgroup default. O accumulator register-resident @ `d1b200b1`. Dest @ `48c56ef` pins `RDNA_ATTN` from API-server `VLLM_USE_RDNA2_FA`. Dest @ `bfd5286d` tile skip / dest **#29** @ `d94e2209` `cu_query_lens` split decode stay extras. Do not dump. Persist O stay extras. | | `csrc/rocm/gdn_decode_rdna2.cu` | `attention/gdn_scan` | 16 fp32 VGPR/thread. SSM `fp16` or `fp32` @ `02adbfd4`. `NULL_BLOCK_ID=0` is a vLLM sentinel — zoo is “invalid slot → zero out”. | | `csrc/rocm/gdn_prefill_*_rdna2.cu` | `attention/gdn_scan` | `o` uses `i_t_local`. Prefill HIP opt-in @ `cd1231fd`. Dispatch still misses dtype. | -| `csrc/rocm/q_gemm_rdna2.cu` + `q_gemm_rdna2_prefill.cu` + `qdq_4_rdna2.cuh` | `gemm/w4a16_fdot2` | One family (`use_v2_format`). ConfigA for `M>256`. Dest deleted `q_gemm_rdna2_awq_prefill.cu` @ `1046782`. ZP still scale-baked. K-split still unaligned. | -| `csrc/rocm/moe_q_gemm_rdna2.cu` | `moe/routed` | Reuses W4 helpers. Dest @ `e1315629` dequant/eight-row stay extras. Do not dump. | +| `csrc/rocm/q_gemm_rdna2.cu` + `q_gemm_rdna2_prefill.cu` + `qdq_4_rdna2.cuh` | `gemm/w4a16_fdot2` | One family (`use_v2_format`). ConfigA for `M>256`. Dest deleted `q_gemm_rdna2_awq_prefill.cu` @ `1046782`. ZP still scale-baked. Dest @ `3a0786ea` K_STEP-aligns unusable POT splits. | +| `csrc/rocm/w4a8_sdot4_rdna2.{cu,cuh}` | — | Dest @ `3a0786ea`. Opt-in `VLLM_RDNA2_W4A8_SDOT4`. gfx1030 prefill `M≥33`. Internal W4A16 fallback. Stay extras. Not a second W4 zoo family. Do not dump. | +| `csrc/rocm/moe_w4a8_rdna2.cu` | — | Dest @ `3a0786ea`. Opt-in. Hard-off under resident MoE. Dest @ `9cad8792` packed-fp16 CAS default. Stay extras. Do not dump. | +| `csrc/rocm/moe_q_gemm_rdna2.cu` + `moe_accum_rdna2.cuh` | `moe/routed` | Reuses W4 helpers. Dest @ `e1315629` dequant/eight-row stay extras. Dest @ `9cad8792` packed-fp16 CAS default; `VLLM_RDNA2_MOE_FP32_ACCUM` opt-in. Do not dump. | | `csrc/rocm/moe_resident_decode.cu` | `moe/routed` (watch) | Dest @ `e1315629`. ATen HIP. `VLLM_RDNA_MOE_RESIDENT*` default **off**. Stay extras. Do not dump. | -| `csrc/rocm/exl3_dot2_{dense,moe,dequant,hadamard}.*` | `gemm/exl3_3inst` | `LDS_PAD=8`. Dest @ `3d6df9ed` widened extras compile guard for docker multi-arch. Zoo still one arch per fatbin. Produce stays `-cb 3inst` outside. Do not dump. | -| `csrc/rocm/causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA, wave32. FIR on pre-shift then shift. | +| `csrc/rocm/exl3_dot2_{dense,moe,dequant,hadamard}.*` | `gemm/exl3_3inst` | `LDS_PAD=8`. Dest @ `3d6df9ed` widened extras compile guard for docker multi-arch. Dest **#32** @ `a5059cf0` mul1 / K=1..8 + MoE n-tile staging bound. Dest @ `7be99d84` fused `exl3_project_rdna2`. Zoo still one arch per fatbin. Produce stays `-cb 3inst` outside. Do not dump. | +| `vllm/.../quantization/exl3.py` | — | Dest @ `aaa374ad` / `30b0bd4e` / `b13effcb`. v0.30 loader + dest EXL3 MTP `mtp.fc`. Stay extras. Python. Do not dump. | +| `scripts/serve_rdna.sh` + `scripts/recipes/*.env` | — | Dest @ `bc5fbee5`. One serve entry + recipes. Stay extras. Serve scripts. Do not copy tok/s. | +| `vllm/.../layers/layernorm.py` | — | Dest @ `74315f4c`. `gemma_rms_norm_fake` contiguous empty. Stay extras. Inductor fake. Do not dump. | +| `tunableop/rocblas-f30bb442e9b5/` + `scripts/serve_gfx1030_exl3_27b.sh` | — | Dest @ `ac9dfd8f`. Canonical TunableOp rows + EXL3 27B launcher. Stay extras. Serve/ops. Do not copy tok/s. | +| `vllm/.../quantization/exl3_moe.py` | — | Dest @ `a50520a3`. `Exl3MoEMethod` for RoutedExperts. Stay extras. Python. Do not dump. | +| `vllm/.../quantization/exl3_ple.py` | — | Dest @ `c7bc9381`. `Exl3NgramTable` gather. Inert until PLE offload rewrite. Stay extras. Python. Do not dump. | +| `csrc/rocm/causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA, wave32. FIR on pre-shift then shift. Dest @ `b13effcb` gates HIP when `state_len != width-1`. Stay extras. Do not dump. | | `csrc/rocm/indexer_paged_mqa_rdna2.cu` | `attention/qsa_indexer` | DeepSeek-class indexer. | | `csrc/rocm/qsa_rdna2.cu` | `attention/qsa_indexer` (watch) | `VLLM_RDNA_QSA_HIP` default **off**. | | `csrc/rocm/sparse_mla_rdna2.cu` | `attention/dsa_nope` | Sparse MLA class. | -| `csrc/rocm/rdna_allreduce.{cu,cuh}` | `comm/pcie` | Uncached+push Later. Dest @ `3b59ee16`: flags in uncached VRAM + T44b wedge. Occupancy pin closed. INT8/Q8 wire; no Finegrained. Do not dump. | +| `csrc/rocm/rdna_allreduce.{cu,cuh}` | `comm/pcie` | Uncached+push Later. Dest @ `3b59ee16`: flags in uncached VRAM + T44b wedge. Dest @ `22d6e346` two-shot; dest @ `101a16c8` two-shot **off** (library `MAX_KB` **64**). Zoo **512**. Do not dump. | +| `vllm/model_executor/layers/hippihx_v1.py` | — | Dest @ `a3f7e5da` (extras **#27**). Opt-in `VLLM_HIPPIHX`. ctypes + `hippihx_v1_load`. Plans not ready. Stay extras. Do not dump. | +| `vllm/model_executor/kernels/linear/mixed_precision/rdna2_w4a16.py` | — | Dest **#30** @ `3fb9d43a`. Opt-in `VLLM_RDNA2_W4A16_RUNTIME_DISPATCH` (`torch.ops.vllm.rdna2_w4a16_gemm`). Default **off**. Python only. Stay extras. Do not dump. | | `csrc/rocm/{mrope,hc,ple_short_conv,rdna_fused_glue}_rdna2.cu` | — | Product HIP. Stay extras. | | `later/glm53-…` `glm5_*` | `kda_scan` / `dsa_nope` / `qsa_indexer` | Later. Drop `glm5_` prefix. | @@ -110,8 +147,9 @@ extras consume. journals FPP13 16k c=8 green via serve arenas — still extras. 2. Tile README LDS / `__launch_bounds__` / wave / DOT unit are filled. 3. hippihx ships one HIP entry extras can bind. **Started:** V1 - `plan`/`run`. `run` is `NOT_READY`. extras @ `30632b2fa323` has no - `torch.ops.hippihx.*`. + `plan`/`run`. `run` is `NOT_READY`. extras @ `bc5fbee5d0a1` has opt-in + `VLLM_HIPPIHX` ctypes consume (dest **#27**). Still no + `torch.ops.hippihx.*`. Plans not ready. 4. extras is rewired **in extras**. The extras copy is then deleted. Until then: observe, lock numbers, keep stubs. Tracker: @@ -119,14 +157,14 @@ Until then: observe, lock numbers, keep stubs. Tracker: ## Dest extras defects -Dest tip `30632b2fa323`. Live dest bugs / rolled-back paths — hippihx +Dest tip `bc5fbee5d0a1`. Live dest bugs / rolled-back paths — hippihx must not reproduce them. -| Defect | Where | Status @ `30632b2` | Zoo lock | +| Defect | Where | Status @ `bc5fbee5` | Zoo lock | |---|---|---|---| | `llvm.amdgcn.fdot2.bf16.bf16` ISel abort | Hybrid W4 gfx10 + bf16. Dest HIP has **no** `fdot2.bf16`. | **Serve-mitigated** (`59237b3`). | Never `fdot2.bf16`. DOT + GDN HIP = **fp16 act**. | | Scale-baked W4 zero-point | `qdq_4_rdna2.cuh` `prep_zero_scale_fp16`: `0xE400 \| zero` then `scale * (-1024 - zero)` in `half`. | **Still live** | Integer `q - zero`, then `* scale`. | -| Unaligned prefill K-split | `compute_split_k`: K=640 → 16×40. Kernel `K_STEP=32`. Cap at 8 and `k_per_split % K_STEP == 0` **not landed**. | **Still live** | Equal `k_per_split`, multiple of 32. Refuse 40-wide. | +| Unaligned prefill K-split | `compute_split_k`: K=640 → 16×40. Kernel `K_STEP=32`. | **Dest-repaired** (`3a0786ea`). Legacy POT kept when `k_per_split % K_STEP == 0`; else K_STEP-aligned enumeration. | Equal `k_per_split`, multiple of 32. Refuse 40-wide. | | Prefill ConfigH `K_STEP=64` | Garbage for `M>256`. | **Dest-reverted** (`7ac98a26`) | Keep ConfigA (`K_STEP=32`). | | GDN HIP selected on BF16 | `_gdn_prefill_dispatch_available()` checks GPU + symbols, **not dtype**. | **Still live** | `plan` / V1 refuse bf16 on `attention.gdn_scan`. Prefill HIP opt-in. | | GDN decode HIP skipped on fp16 SSM state | Dispatch required `ssm_state.dtype == float32`. | **Dest-fixed** (`02adbfd4`) | Act **fp16**. SSM **fp16 or fp32**. Refuse bf16 act. | @@ -147,6 +185,7 @@ must not reproduce them. | gfx1030 wvSplitK decode (`n<=5`) | `wvSplitK_hf_big_` device `assert(false)` under cudagraph capture | **Dest-reverted** (`5c4ab989`; was `c350fa218`). Decode stays `gemv_f16_rdna2` `M<=8`. | **No zoo tile.** Do not grow `gemv_f16`. | | QSA prefix cache always-zero | `CircularBufferManager.find_longest_cache_hit` returned 0 at empty group ring | **Dest-fixed (serve)** (`e45dd5cb`) | Stay extras. Empty ring at group boundary is valid. | | Mamba spec-decode batch-row tables | captured `data_ptrs` indexed by batch row under deferred postprocess | **Dest-fixed (serve)** (`741e5bc3`). V1 unchanged. | Stay extras. Persistent per-request-slot tables; index by `req_idx`. | +| EXL3 MoE trellis staging over-read | `exl3_dot2_moe.cu` staged 64 n-tiles when `N<1024` | **Dest-fixed** (`a5059cf0`) | Stage only the block's live n-tiles. | Serve notes (not zoo): GDN prefill HIP opt-in (`cd1231fd`); RDNA_ATTN spec/MTP verify default-off; `eager_break_during_capture` stays extras; @@ -183,6 +222,8 @@ This review is documentation, not a kernel migrate. | dest Flash-Next / Hybrid W4 / recipe ports | Unique: **Aron Hsiao** `5765f57b4c41` / `c05af408775f` / `22bb2e8d06f3` | Their unique commits if dest-locked. Dest Blivion follow-ups stay Blivion | | PR #2 `glm5_kda_*` / `glm5_dsa_*` | BlivionIaG | BlivionIaG; rename off `glm5_` in a follow-up hippihx commit | | PR #3 a17t unique W4 (`d53572644`, later Simon Siebert) | **Not taken** | Closed unmerged. If dest ever locks that family, pick **their** commits | -| Explore PRs **#9/#10** sdot | **Not taken** | Not dest | -| extras dest-presence serve (**#13** T44b / **#14** / **#15** / **#17** / **#19** / **#20** / dest-integrated **#21** / mamba `741e5bc3` / closed **#12**) | observe | Python/Triton/serve + dest-landed ATen HIP. No HIP migrate. Do not pick Cursor/George/Codex/Karl/Waldecir/a17t rewrites. Unique Aron Hsiao `rdna_ar` still pickable. | -| extras not dest (**#16** closed / **#18** / **#22** / **#23** / **#24**) | **Not taken** | GPTQ `BLOCK_KN_SIZE` 256. Two-shot `rdna_ar`. PCIe P2P KV. RAM KV offload. | +| Explore PRs **#9/#10** sdot | **Not taken** | **#9** closed-unmerged (superseded by dest W4A8 @ `3a0786ea`). **#10** still not dest. | +| dest W4A8 sdot4 (`w4a8_sdot4_rdna2.*`, `moe_w4a8_rdna2.cu`) | dest @ `3a0786ea` BlivionIaG | Stay extras. Opt-in. Not a second W4 zoo family. Do not dump. | +| dest MoE epilogue (`moe_accum_rdna2.cuh`) | dest @ `9cad8792` BlivionIaG | Stay extras. Packed-fp16 CAS default. `VLLM_RDNA2_MOE_FP32_ACCUM` opt-in. Do not dump. | +| extras dest-presence serve (dest-landed extras PRs, dest **#29** split decode, dest row-gate `e0112c55`, dest W4A8 `3a0786ea`, dest **#30** `3fb9d43a`, dest MoE epilogue `9cad8792`, dest **#32** `a5059cf0`, dest EXL3 loader `30b0bd4e`, dest TunableOp `ac9dfd8f`, dest EXL3 MTP + `serve_rdna` `bc5fbee5`, mamba `741e5bc3`) | observe | Python/Triton/serve + dest-landed ATen HIP. No HIP migrate. Do not pick Cursor/George/Codex/Karl/Waldecir/a17t/Claude rewrites. Unique Aron Hsiao `rdna_ar` still pickable. | +| extras not dest (**#16** closed / **#18** / **#23** / **#25** / **#31** / **#33** / **#34**) | **Not taken** | GPTQ `BLOCK_KN_SIZE` 256. PCIe P2P KV. rdna_ar retained-output. W4 exact-dequant explore. FA decode scores. FA D=128 prefill. | diff --git a/docs/CONSUME.md b/docs/CONSUME.md index ae6ae7b..558c18e 100644 --- a/docs/CONSUME.md +++ b/docs/CONSUME.md @@ -79,9 +79,11 @@ fa_fdot2.run(binding) Register **one** `torch.ops.hippihx.` that calls `hippihx_v1_run`. No Triton→HIP double-fire. Bind keys on **arch + wave**, not GFX name alone. extras V1 consume is the HIP fatbin. Dest extras tip -`30632b2fa323` still has no `torch.ops.hippihx` bind. FlyDSL extras -consume waits on `FLYDSL_V1_CONSUME` (graph-safe JIT) — see -[`FLYDSL.md`](FLYDSL.md). +`bc5fbee5d0a1` still has no `torch.ops.hippihx` bind. Dest extras **#27** +(`a3f7e5da`) is opt-in `VLLM_HIPPIHX` ctypes consume via +`hippihx.v1_ctypes` + `hippihx_v1_load` (default **off**; every plan +still not ready). FlyDSL extras consume waits on `FLYDSL_V1_CONSUME` +(graph-safe JIT) — see [`FLYDSL.md`](FLYDSL.md). ## Host contracts already here diff --git a/docs/EXTRAS.md b/docs/EXTRAS.md index eafeb57..b0e0128 100644 --- a/docs/EXTRAS.md +++ b/docs/EXTRAS.md @@ -7,22 +7,28 @@ one V1 op. Review: [`BACKPORT.md`](BACKPORT.md). ## Unvalidated extras inventory Snapshot of [`opengfx1030/vllm-rdna`](https://github.com/opengfx1030/vllm-rdna) -`rdna_extras` @ `30632b2fa323` (2026-09-27 13:05 UTC). Dest default +`rdna_extras` @ `bc5fbee5d0a1` (2026-10-01 13:08 UTC). Dest default branch is **`rdna_extras`**. `main` is upstream vLLM `c00091e02670`. Merged extras [PR #1](https://github.com/opengfx1030/vllm-rdna/pull/1) -squash is `a4060647cfbb`. Open **#2**, **#18**, **#22**, **#24**. Draft **#23**. -Closed **#3/#16** unmerged; **#21** dest-integrated unmerged; **#12** -superseded by **#15**. Merged **#13**, **#14**, **#15**, **#17**, **#19**, -**#20**. -**No** `torch.ops.hippihx.*`. +squash is `a4060647cfbb`. Open **#2**, **#18**. Draft **#10**, **#23**, +**#25**, **#31**, **#33**, **#34**. Closed **#3/#9/#16/#30/#32** unmerged +(**#9** superseded by dest W4A8; **#30** dest-landed independently @ +`3fb9d43a`; **#32** dest-landed independently @ `a5059cf0`); **#21** +dest-integrated unmerged; **#12** superseded by **#15**. Merged **#13**, +**#14**, **#15**, **#17**, **#19**, **#20**, **#22**, **#24**, **#26**, +**#27**, **#28**, **#29**. Dest **#30** @ `3fb9d43a` / dest **#32** @ +`a5059cf0` dest-presence. Dest tip `bc5fbee5` is serve recipes + EXL3 +MTP (not a new extras PR). **No** `torch.ops.hippihx.*`. Dest **#27** is +opt-in `VLLM_HIPPIHX` ctypes consume (default **off**; plans not ready). **Unvalidated.** Not dest. Not silicon-signed. No tok/s. hippihx still ships stubs. Observed SHAs and actions live in [`BACKPORT.md`](BACKPORT.md). Live dest -bugs (do not copy): W4 scale-baked ZP, unaligned K-split, GDN HIP-on-BF16, -ConfigH, GDN batched-decode n≥8. Dest retraces Flash-Next -FULL_AND_PIECEWISE c=8 to probe artifacts (`609c9c0d`). +bugs (do not copy): W4 scale-baked ZP, GDN HIP-on-BF16, +ConfigH, GDN batched-decode n≥8. Dest @ `3a0786ea` K_STEP-aligns +unusable POT K-splits. Dest retraces Flash-Next FULL_AND_PIECEWISE c=8 +to probe artifacts (`609c9c0d`). Status: **extras** = live on dest tip · **Later** = side branch · **skip** = do not take · **stub** = hippihx contract only. @@ -31,30 +37,44 @@ Status: **extras** = live on dest tip · **Later** = side branch · | Kernel / op | extras | hippihx tile | Notes | |---|---|---|---| -| FA paged decode / prefill / split-K / short / GQA | `fa_rdna2.cu` | `attention/fa_fdot2` | `fdot2`. Occupancy pin closed. GQA-subgroup default. O register-resident @ `d1b200b1`. Persist O stay extras. **unvalidated** | +| FA paged decode / prefill / split-K / short / GQA | `fa_rdna2.cu` | `attention/fa_fdot2` | `fdot2`. Occupancy pin closed. GQA-subgroup default. O register-resident @ `d1b200b1`. Dest **#28** @ `bfd5286d` tile skip / in-place out / GQA softmax stay extras. Do not dump. Persist O stay extras. **unvalidated** | | FA INT8 KV writer | `reshape_and_cache_int8_rdna2` | `attention/fa_fdot2` | INT8 cache layout. **unvalidated** | | FA fp16 flash KV writer | `reshape_and_cache_flash_rdna2` | `attention/fa_fdot2` | Non-native KV. `__launch_bounds__(128, 4)`. **unvalidated** | | W4A16 dense decode | `q_gemm_rdna2.cu` | `gemm/w4a16_fdot2` | GPTQ + AWQ = pack/zeros, one GEMM. Dest ZP is scale-baked `half` — zoo uses integer `q-zero` then scale. **unvalidated** | -| W4A16 prefill | `q_gemm_rdna2_prefill.cu` | `gemm/w4a16_fdot2` | Unified GPTQ+AWQ. ConfigA for `M>256`. Dest reverted ConfigH. Dest can pick K=640→16×40; zoo requires equal ×32. **unvalidated** | +| W4A16 prefill | `q_gemm_rdna2_prefill.cu` | `gemm/w4a16_fdot2` | Unified GPTQ+AWQ. ConfigA for `M>256`. Dest reverted ConfigH. Dest @ `3a0786ea` K_STEP-aligns unusable POT splits. Zoo still refuse 40-wide. **unvalidated** | +| W4A8 sdot4 dense prefill | `w4a8_sdot4_rdna2.{cu,cuh}` (dest @ `3a0786ea`) | — | Stay extras. Opt-in `VLLM_RDNA2_W4A8_SDOT4`. gfx1030 `M≥33`. Internal W4A16 fallback. Not a second W4 zoo family. Do not dump. **unvalidated** | +| MoE W4A8 sdot4 | `moe_w4a8_rdna2.cu` (dest @ `3a0786ea`) | — | Stay extras. Opt-in. Hard-off under resident MoE. Dest @ `9cad8792` packed-fp16 CAS default. Do not dump. **unvalidated** | | W4A16 AWQ high-M prefill | ~~`q_gemm_rdna2_awq_prefill.cu`~~ | `gemm/w4a16_fdot2` | **Dest-deleted** @ `1046782`. Do not reintroduce. | -| W4A16 MoE | `moe_q_gemm_rdna2.cu` | `moe/routed` | Same W4 family. Dest @ `e1315629` dequant/eight-row stay extras. moe_align prealloc is extras. **unvalidated** | -| EXL3 dense / MoE / dequant / Hadamard / trellis decode | `exl3_dot2_*.cu` | `gemm/exl3_3inst` | Consume `-cb 3inst`. Produce outside. UNC-26. **unvalidated** | +| W4A16 MoE | `moe_q_gemm_rdna2.cu` + `moe_accum_rdna2.cuh` | `moe/routed` | Same W4 family. Dest @ `e1315629` dequant/eight-row stay extras. Dest @ `9cad8792` packed-fp16 CAS default; `VLLM_RDNA2_MOE_FP32_ACCUM` opt-in. moe_align prealloc is extras. **unvalidated** | +| EXL3 dense / MoE / dequant / Hadamard / trellis decode | `exl3_dot2_*.cu` | `gemm/exl3_3inst` | Consume `-cb 3inst`. Produce outside. Dest **#32** @ `a5059cf0` mul1 / K=1..8 + MoE n-tile bound. Dest @ `7be99d84` fused `exl3_project_rdna2`. UNC-26. **unvalidated** | | GDN packed decode | `gdn_decode_rdna2.cu` | `attention/gdn_scan` | Dest @ `02adbfd4`: SSM **fp16 or fp32**. Dest @ `388a61b6`: no one-shot `zero_()` wipe. **fp16 act only.** **unvalidated** | | GDN prefill chain | `gdn_prefill_*_rdna2.cu` | `attention/gdn_scan` | **Opt-in** (`VLLM_GDN_HIP_PREFILL=1`; default Triton/FLA @ `cd1231fd`). `o` varlen `i_t_local` dest-fixed. Dispatch still misses dtype. **unvalidated** | -| causal_conv1d update + fwd | `causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA. FIR on pre-shift then shift. BF16: fp32 mul, **not** `fdot2.bf16`. **unvalidated** | +| causal_conv1d update + fwd | `causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA. FIR on pre-shift then shift. Dest @ `b13effcb` gates HIP when `state_len != width-1`. BF16: fp32 mul, **not** `fdot2.bf16`. **unvalidated** | | Paged MQA indexer | `indexer_paged_mqa_rdna2.cu` | `attention/qsa_indexer` | gfx1030 BF16 6h×256: **4 warps**. **unvalidated** | | Flash-Next QSA store/compress/MQA | `qsa_rdna2.cu` | `attention/qsa_indexer` (watch) | `VLLM_RDNA_QSA_HIP` default **off**. Eager Triton serves (`8cf0dedb`). Dest @ `e45dd5cb`: empty QSA ring prefix hits. Dest @ `f3dd65fa`: live-context prefill scoring bound (Triton). **unvalidated** | | Sparse MLA decode / prefill | `sparse_mla_rdna2.cu` | `attention/dsa_nope` | **unvalidated** | | M-RoPE / Flash-Next HC / PLE / fused glue | `mrope_rdna2.cu`, `hc_rdna2.cu`, `ple_short_conv_rdna2.cu`, `rdna_fused_glue.cu` | — | Product HIP. HC/PLE default off (S6 revert `9c9509b3`); wrappers @ `d0d577f1`; isolated HC @ `8960a3bc`; `_contig()` cache @ `50120e13` (gate still off). Stay extras. **unvalidated** | | W8A16 / FP8 / MXFP4 / gfx1100 WMMA / skinny GEMM / RMSNorm | `moe_w8a16*.cu`, `mxfp4_dot2_*.cu`, `q_gemm_rdna3_wmma.cu`, `skinny_gemms*.cu`, `layernorm.cu` | — | No tile. WMMA is Later overlay. Do not reintroduce leapdragon `gemv_f16`. **unvalidated** | | GLM-5.3 KDA / DSA | `glm5_*.cu` (PR **#2**) | `kda_scan` / `dsa_nope` / `qsa_indexer` | Later. Drop `glm5_` name. **unvalidated** | -| leapdragon push AR | `rdna_allreduce.{cu,cuh}` (merged PR **#1**) | `comm/pcie` | Dest extras, default **off**. Occupancy pin closed. **unvalidated** | +| leapdragon push AR | `rdna_allreduce.{cu,cuh}` (merged PR **#1**; dest **#22** @ `22d6e346`) | `comm/pcie` | Dest extras, default **off**. Dest two-shot then dest-off @ `101a16c8` (library `MAX_KB` **64**). Zoo **512**. Do not dump. **unvalidated** | | a17t extra AWQ GEMM / GEMV | `awq_gemm_rdna2.cu` etc. (closed PR **#3**) | — | **skip** — closed unmerged, second W4 family | | GPTQ exllama `BLOCK_KN_SIZE` 256 | `q_gemm.cu` (PR **#18**) | — | **skip** — not dest DOT W4 | -| two-shot `rdna_ar` | `rdna_allreduce.{cu,cuh}` (PR **#22**) | — | **skip** — not dest. Library `VLLM_RDNA_AR` still `"0"`. Zoo `AR_MAX_KB` **512**. Do not dump | | PCIe P2P KV disagg | `PcieP2pConnector` (draft PR **#23**) | — | **skip** — not dest. HIP IPC SDMA. Zoo `comm.pcie` is AR, not KV. Do not dump | -| RAM KV offload | offloading connector (PR **#24**) | — | **skip** — not dest. George/Codex. Serve only. No HIP. Do not pick | -| Explore W4A8 sdot4 / W4A4 sdot8 | extras PRs **#9/#10** | — | **skip** — not dest | +| RAM KV offload | offloading connector (dest **#24** @ `cd38a1d3`) | — | Stay extras. George/Codex. Serve only. No HIP. Do not pick | +| rdna_ar retained-output | draft PR **#25** | — | **skip** — not dest. George/Codex. Do not dump. Do not pick | +| HIP MoE skinny `MAX_M=16` | dest PR **#26** @ `7434efee` | — | Stay extras. Opt-in `VLLM_ROCM_MOE_SKINNY_MAX_M`. Default still **8**. Do not dump. Do not pick | +| hippihx V1 consume | `hippihx_v1.py` (dest **#27** @ `a3f7e5da`) | — | Stay extras. Opt-in `VLLM_HIPPIHX`. ctypes + `hippihx_v1_load`. Plans not ready. Still no `torch.ops.hippihx`. Do not dump. Do not pick Claude | +| FA-RDNA2 tile skip / in-place / GQA softmax | dest PR **#28** @ `bfd5286d` | `attention/fa_fdot2` | Stay extras. Foreign Claude. Do not dump. Do not pick. **unvalidated** | +| FA spec-decode split | dest PR **#29** @ `d94e2209` / tip `e0112c55` | `attention/fa_fdot2` | Stay extras. Dest rework of foreign Claude (`cu_query_lens` decode + row gate ≤256). Supersedes dest @ `83e6af80` per-position verify-decode. Do not dump. Do not pick Claude. **unvalidated** | +| W4 `torch.compile` M-dispatch | dest PR **#30** @ `3fb9d43a` | — | Stay extras. Opt-in `VLLM_RDNA2_W4A16_RUNTIME_DISPATCH`. Default **off**. Python only. GitHub PR closed-unmerged. Do not dump. Do not copy tok/s. | +| W4 exact-dequant explore | draft PR **#31** | — | **skip** — explore, not dest. Claude. Opt-in `VLLM_RDNA2_W4A16_EXACT_DEQUANT`. Zoo already locks integer `q-zero` then scale. Do not dump | +| EXL3 mul1 decode / K=1..8 trellis | dest PR **#32** @ `a5059cf0` | `gemm/exl3_3inst` | Stay extras. Unsigned mul1 + K=1..8 + dest-fixed MoE staging bound. GitHub PR closed-unmerged. Do not dump. | +| EXL3 v0.30 loader / fused project / MoE / PLE | dest @ `30b0bd4e` (`exl3.py`, `exl3_moe.py`, `exl3_ple.py`, `exl3_project_rdna2`) | `gemm/exl3_3inst` | Stay extras. Codebook markers, suh grouping, TP slicing, `fold_weight` lm_head only, `Exl3MoEMethod`, inert `Exl3NgramTable`. Unique Blivion. Do not dump. | +| TunableOp unify / EXL3 27B rows / gemma_rms fake | dest @ `ac9dfd8f` (`tunableop/rocblas-f30bb442e9b5`, `gemma_rms_norm_fake`) | — | Stay extras. Serve/ops + inductor fake. Unique Blivion. Do not dump. Do not copy tok/s. | +| EXL3 MTP draft + `serve_rdna` recipes | dest tip `bc5fbee5` (`exl3.py` `mtp.fc`, `serve_rdna.sh`, HIP conv `state_len` gate) | — | Stay extras. Dest serves EXL3 quantized MTP. No zoo MTP tile. Unique Blivion. Do not dump. Do not copy tok/s. | +| FA decode scores / 4 barriers | draft PR **#33** | — | **skip** — not dest. Do not dump `fa_rdna2` | +| FA D=128 prefill register-O GQA | draft PR **#34** | — | **skip** — not dest. Do not dump `fa_rdna2` | +| Explore W4A8 sdot4 / W4A4 sdot8 | extras PRs **#9/#10** | — | **#9** closed-unmerged (superseded by dest W4A8 @ `3a0786ea`). **#10** **skip** — not dest | | Resident W4A16 MoE skinny decode | `moe_resident_decode.cu` (dest **#17** @ `e1315629`) | `moe/routed` (watch) | Stay extras. Opt-in `VLLM_RDNA_MOE_RESIDENT*`. ATen HIP. Do not dump. Closed **#16** unmerged. **unvalidated** | ### Modes / dispatch @@ -72,11 +92,16 @@ Status: **extras** = live on dest tip · **Later** = side branch · | causal conv HIP | update + fwd | on unless set `0` | `causal_conv` | | Flash-Next HC/QSA/PLE HIP | dest scaffolding | **off**; S6 default-on reverted; wrappers @ `d0d577f1`; isolated HC @ `8960a3bc`; `_contig()` cache @ `50120e13` (same-shape clobber still live) | extras until dest-on | | Hybrid W4A16 gfx10 | extras linear backend | ungated; RDNA2 W4 auto on gfx1030 | not a second W4 family | +| W4A8 sdot4 prefill | dest @ `3a0786ea` int4×int8 `sdot4` | **off** (`VLLM_RDNA2_W4A8_SDOT4` unset = W4A16) | extras; not a second W4 zoo family | | gfx1030 wvSplitK n≤5 | extras dense GEMM (`c350fa218`) | **dest-reverted** (`5c4ab989`); decode stays `gemv_f16_rdna2` `M<=8` | **no tile** | -| V1 FULL_AND_PIECEWISE | dest captures FULL + piecewise | extras runner (`1ff73596`); dest @ `68a635ed` keeps FULL decode graphs; piecewise for mixed/prefill | serve | +| V1 FULL_AND_PIECEWISE | dest captures FULL + piecewise | extras runner (`1ff73596`); dest @ `68a635ed` keeps FULL decode graphs; dest @ `e0112c55` `UNIFORM_BATCH` when split-decode + verify FULL graph | serve | +| FA split decode | dest **#29** decode-first mixed/verify | **on** (`VLLM_FA_RDNA2_SPLIT_DECODE=1`); row gate ≤256 query×head @ `e0112c55` | extras | | Custom AR (vLLM/ROCm) | force custom all-reduce on PCIe | **on** in dest gfx1030 launcher (`4d25a048`); `envs.py` still False; cudagraph-correct @ `849292ec` | serve | -| leapdragon `rdna_ar` | size-gated Uncached+push | **off** (`VLLM_RDNA_AR=0`) | `comm/pcie` Later | +| leapdragon `rdna_ar` | size-gated Uncached+push; dest **#22** two-shot dest-off @ `101a16c8` | **off** (`VLLM_RDNA_AR=0`) | `comm/pcie` Later | | Resident W4A16 MoE | native shuffled layout + skinny GEMV | **off** (dest **#17** @ `e1315629`) | extras | +| hippihx V1 consume | dest **#27** ctypes + `hippihx_v1_load` | **off** (`VLLM_HIPPIHX=0`) | extras; plans not ready | +| W4 compile-dispatch | dest **#30** `torch.ops.vllm.rdna2_w4a16_gemm` | **off** (`VLLM_RDNA2_W4A16_RUNTIME_DISPATCH=0`) | extras; dest keeps default off | +| MoE epilogue accum | dest @ `9cad8792` packed-fp16 CAS | **off** (`VLLM_RDNA2_MOE_FP32_ACCUM=0`); `=1` is fp32 scratch | extras | ### Env (extras-added / extras-used) @@ -103,8 +128,17 @@ stay extras (no D2H under capture in the zoo). | `VLLM_EXL3_FOLDED_CACHE` | unset | folded-weight cache dir | | `VLLM_ROCM_USE_SKINNY_GEMM` | `True` | skinny GEMM | | `VLLM_ROCM_MOE_SKINNY` | `1` | MoE skinny | +| `VLLM_ROCM_MOE_SKINNY_MAX_M` | `8` (dest **#26** @ `7434efee`; opt-in `16`) | sequential HIP MoE row cap | +| `VLLM_HIPPIHX` | `False` (dest **#27** @ `a3f7e5da`; only `"1"` enables) | opt-in hippihx V1 consume | +| `VLLM_RDNA2_W4A8_SDOT4` | unset / `"0"` (dest @ `3a0786ea`; only `"1"` enables) | opt-in W4A8 sdot4 prefill | +| `VLLM_RDNA2_W4A16_RUNTIME_DISPATCH` | `"0"` (dest **#30** @ `3fb9d43a`; only `"1"` enables) | opt-in W4 compile-time M-dispatch via custom op | +| `VLLM_RDNA2_MOE_FP32_ACCUM` | `False` (dest @ `9cad8792`; only `"1"` enables) | opt-in MoE fp32 scratch epilogue; dest default is packed-fp16 CAS | +| `VLLM_HIPPIHX_LIB` / `VLLM_HIPPIHX_CODE_OBJECT` | unset | `libhippihx_v1.so` / `hippihx_.hsaco` | | `VLLM_RDNA_MOE_RESIDENT` / `VLLM_RDNA_MOE_RESIDENT_SKINNY` | `"0"` (dest **#17** @ `e1315629`) | resident W4A16 MoE layout / skinny decode | | `VLLM_FA_RDNA2_GQA_MODE` | `subgroup` | FA GQA-subgroup prefill | +| `VLLM_FA_RDNA2_GQA_DECODE` | `"0"` (dest **#28** @ `bfd5286d`; dest launchers set `1` when FA) | opt-in GQA decode kernel | +| `VLLM_FA_RDNA2_SPLIT_DECODE` | `"1"` (dest **#29** @ `d94e2209`) | decode-first mixed/verify split | +| `VLLM_FA_RDNA2_VERIFY_FULL_GRAPH` | `"1"` (dest **#29** dest-on; PR was opt-in) | `UNIFORM_BATCH` for uniform verify | | `VLLM_RDNA_HC_PREFILL_HIP` / `VLLM_RDNA_QSA_HIP` / `VLLM_RDNA_PLE_CONV_HIP` | `"0"` | Flash-Next product HIP (opt-in) | | `VLLM_RDNA_FUSED_HC` / `VLLM_RDNA_FUSED_SE` | `"1"` (Flash-Next production launcher sets fused HC `0` @ `3bddd3c9`) | fused decode glue | | `VLLM_ROCM_USE_AITER` | `False` | AITER (CDNA; not dest gfx1030) | @@ -116,7 +150,7 @@ stay extras (no D2H under capture in the zoo). | `VLLM_RDNA_AR` | `"0"` (dest extras; merged PR **#1**) | leapdragon push AR (opt-in; communicator gate, not the stale “enabled by default” docstring) | | `VLLM_RDNA_AR_BLOCKS` | auto | AR block cap | | `VLLM_RDNA_AR_PACE` | `0` | AR store pace | -| `VLLM_RDNA_AR_MAX_KB` | dest extras **64** (`3b59ee16`). zoo lock **512** | AR fast-path size cap | +| `VLLM_RDNA_AR_MAX_KB` | dest extras **64** (`101a16c8`). zoo lock **512** | AR fast-path size cap | | `VLLM_USE_BREAKABLE_CUDAGRAPH` | `0` (auto-on in some configs) | capture dispatcher | | `VLLM_LOG_GDN_PTRS` / `VLLM_GDN_DBG` / `VLLM_EXL3_*_DBG` / `VLLM_CONV1D_DEBUG` / `DBG_VLLM_STEP_TIMING` | off | probes | @@ -128,19 +162,19 @@ stay extras (no D2H under capture in the zoo). | Custom all-reduce (vLLM/ROCm) | `VLLM_FORCE_CUSTOM_ALL_REDUCE` | **on** in dest gfx1030 launcher (`4d25a048`); envs.py still False | serve | | AITER custom AR | `VLLM_ROCM_USE_AITER_CUSTOM_AR` | on in envs, AITER itself off | CDNA, not gfx1030 dest | | Quick-reduce / symm-mem AR | `VLLM_ROCM_QUICK_REDUCE_*` / `VLLM_ALLREDUCE_USE_SYMM_MEM` | unset / on | serve | -| leapdragon `rdna_ar` Uncached+push | dest extras (PR **#1** @ `a4060647`; T44b @ `3b59ee16`) | **off** | `comm/pcie` Later. Unique HIP: Aron Hsiao. Do not pick Cursor squash. | +| leapdragon `rdna_ar` Uncached+push | dest extras (PR **#1** @ `a4060647`; T44b @ `3b59ee16`; dest **#22** then dest-off @ `101a16c8`) | **off** | `comm/pcie` Later. Unique HIP: Aron Hsiao. Dest `MAX_KB` **64**. Zoo **512**. Do not pick Cursor squash. | ### Not taken / leave in extras -Stay extras: capture plumbing, product serve, skinny GEMM, ConfigH, -graph mode, vLLM custom AR, `bench_report.py`, dest T44b `rdna_ar` -(still opt-in), dest PR **#13** / PR **#14** / PR **#15** / PR **#17** / -PR **#19** / PR **#20** (observe, do not pick), dest-integrated PR **#21**, -dest FA `RDNA_ATTN` pin, resident MoE ATen HIP, Qwen4Exp MTP (not dest), -ROCm platform init, extras EXL3 docker arch-guard. Skip: a17t PR **#3**, -PR **#18** GPTQ `BLOCK_KN_SIZE` 256, PR **#22** two-shot `rdna_ar` -(Cursor; library `VLLM_RDNA_AR` still `"0"`; zoo `AR_MAX_KB` **512**; -do not dump), draft PR **#23** PCIe P2P KV (Cursor; HIP IPC SDMA; do not -dump), PR **#24** RAM KV offload (George/Codex; no HIP; do not pick), -explore **#9/#10**, closed PR **#12**, closed PR **#16**, -closed **#5/#11**. Produce stays outside hippihx. +Stay extras: serve, product HIP, dest-landed extras PRs, dest **#29** +split decode / dest row-gate `e0112c55` / dest W4A8 `3a0786ea` / dest +**#30** compile-dispatch `3fb9d43a` / dest MoE epilogue `9cad8792` / +dest **#32** EXL3 mul1 `a5059cf0` / dest EXL3 loader `30b0bd4e` / dest +TunableOp `ac9dfd8f` / dest EXL3 MTP + `serve_rdna` `bc5fbee5` +(observe, do not pick). Skip: a17t +PR **#3**, PR **#18** GPTQ `BLOCK_KN_SIZE` 256, draft PR **#23** PCIe +P2P KV, draft PR **#25** rdna_ar retained-output, draft PR **#31** W4 +exact-dequant explore, draft PR **#33** FA decode scores, draft PR +**#34** FA D=128 prefill, closed **#9** (superseded), closed **#30** / +**#32** (dest-landed independently), explore **#10**, closed PR **#12**, +closed PR **#16**, closed **#5/#11**. Produce stays outside hippihx. diff --git a/tests/test_no_b12x_cuda.py b/tests/test_no_b12x_cuda.py index 638e09a..856860c 100644 --- a/tests/test_no_b12x_cuda.py +++ b/tests/test_no_b12x_cuda.py @@ -112,7 +112,24 @@ def test_readme_unvalidated_inventory() -> None: assert "e131562" in extras assert "48c56ef" in extras assert "68a635" in extras - assert "30632b2" in extras + assert "bfd5286d" in extras + assert "83e6af80" in extras + assert "e0112c55" in extras + assert "3a0786ea" in extras + assert "3fb9d43a" in extras + assert "9cad8792" in extras + assert "a5059cf0" in extras + assert "30b0bd4e" in extras + assert "ac9dfd8f" in extras + assert "bc5fbee5" in extras + assert "exl3_project_rdna2" in extras + assert "gemma_rms_norm_fake" in extras + assert "serve_rdna" in extras + assert "a3f7e5da" in extras + assert "VLLM_HIPPIHX" in extras + assert "VLLM_RDNA2_W4A8_SDOT4" in extras + assert "VLLM_RDNA2_W4A16_RUNTIME_DISPATCH" in extras + assert "VLLM_RDNA2_MOE_FP32_ACCUM" in extras assert "e45dd5c" in extras assert "d1b200" in extras assert "1046782" in extras @@ -148,6 +165,26 @@ def test_readme_unvalidated_inventory() -> None: assert "2a5e893" in bp assert "700753d9" in bp assert "30632b2" in bp + assert "cd38a1d3" in bp + assert "22d6e346" in bp + assert "a3f7e5da" in bp + assert "7434efee" in bp + assert "101a16c8" in bp + assert "330b42ab" in bp + assert "bfd5286d" in bp + assert "83e6af80" in bp + assert "e0112c55" in bp + assert "3a0786ea" in bp + assert "3fb9d43a" in bp + assert "9cad8792" in bp + assert "a5059cf0" in bp + assert "30b0bd4e" in bp + assert "ac9dfd8f" in bp + assert "bc5fbee5" in bp + assert "exl3_project_rdna2" in bp + assert "gemma_rms_norm_fake" in bp + assert "serve_rdna" in bp + assert "d94e2209" in bp assert "dbb1e77" in bp assert "e45dd5c" in bp assert "d1b200" in bp @@ -186,13 +223,39 @@ def test_readme_unvalidated_inventory() -> None: assert "PR **#22**" in extras or "PR **#22**" in bp assert "PR **#23**" in extras or "PR **#23**" in bp assert "PR **#24**" in extras or "PR **#24**" in bp + assert "PR **#25**" in extras or "PR **#25**" in bp + assert "PR **#26**" in extras or "PR **#26**" in bp + assert "PR **#27**" in extras or "PR **#27**" in bp + assert "PR **#28**" in extras or "PR **#28**" in bp + assert "PR **#29**" in extras or "PR **#29**" in bp + assert "PR **#30**" in extras or "PR **#30**" in bp + assert "PR **#31**" in extras or "PR **#31**" in bp + assert "PR **#32**" in extras or "PR **#32**" in bp + assert "PR **#33**" in extras or "PR **#33**" in bp + assert "PR **#34**" in extras or "PR **#34**" in bp assert "VLLM_RDNA_MOE_RESIDENT" in extras assert "moe_resident" in extras contrib = (ROOT / "CONTRIBUTING.md").read_text(encoding="utf-8") - assert "30632b2" in contrib + assert "bc5fbee5" in contrib assert "a4060647" in contrib # CONTRIBUTING is tip-only. Historical dest SHAs live in BACKPORT. for token in ( + "ac9dfd8f", + "30b0bd4e", + "a5059cf0", + "9cad8792", + "3fb9d43a", + "3a0786ea", + "e0112c55", + "83e6af80", + "d94e2209", + "bfd5286d", + "330b42ab", + "30632b2", + "bcdaaddc", + "22d6e346", + "101a16c8", + "a3f7e5da", "700753d9", "2a5e893", "68a635", diff --git a/tiles/attention/fa_fdot2/README.md b/tiles/attention/fa_fdot2/README.md index 36b2f82..2c560ff 100644 --- a/tiles/attention/fa_fdot2/README.md +++ b/tiles/attention/fa_fdot2/README.md @@ -37,7 +37,10 @@ body yet — when you do, cherry-pick BlivionIaG, do not re-author Dest extras @ `48c56ef` pins `AttentionConfig.backend = RDNA_ATTN` from the API-server `VLLM_USE_RDNA2_FA` env so workers actually select -FA-RDNA2. Stay extras. No HIP. +FA-RDNA2. Dest extras **#28** @ `bfd5286d` and dest **#29** @ `d94e2209` / +`e0112c55` (`cu_query_lens` split decode + row gate) stay extras. +Dest tip `bc5fbee5` does not change FA HIP. Do not dump. Do not pick +Claude. No tok/s. ## V1 rev 4 contract (pinned) diff --git a/tiles/gemm/exl3_3inst/README.md b/tiles/gemm/exl3_3inst/README.md index f33ac02..8198907 100644 --- a/tiles/gemm/exl3_3inst/README.md +++ b/tiles/gemm/exl3_3inst/README.md @@ -20,7 +20,13 @@ per fatbin. gfx1150 / gfx12xx are not dest. Do not dump the `.cu`. | Produce | **outside** hippihx — do not add a `produce/` or `3inst/` packer dir here | | Consume ABI | `HIPPIHX_V1_OP_GEMM_EXL3_3INST` in `include/hippihx/v1.h` — plan/run stub until body migrate | -`mcg` / `mul1` codebooks are compile slots, not extra zoo roots. UNC-26 is +`mcg` / `mul1` codebooks are compile slots, not extra zoo roots. Dest +extras **#32** @ `a5059cf0` dest-lands unsigned mul1 + K=1..8 dispatch +and dest-fixes MoE n-tile staging (GitHub closed-unmerged). Dest @ +`30b0bd4e` dest-lands the v0.30 loader, fused `exl3_project_rdna2`, +`Exl3MoEMethod`, and inert `Exl3NgramTable` PLE gather. Dest @ +`b13effcb` serves the EXL3 quantized MTP draft. Dest tip `bc5fbee5` +is serve recipes. Stay extras. Do not dump. UNC-26 is still In Progress on extras — do not copy bodies until produce/consume and capture buffers stop moving. Port and later perf commits stay **BlivionIaG** `` as Author **and** Committer. Do not diff --git a/tiles/sequence/causal_conv/README.md b/tiles/sequence/causal_conv/README.md index 87ac376..ec261bd 100644 --- a/tiles/sequence/causal_conv/README.md +++ b/tiles/sequence/causal_conv/README.md @@ -14,7 +14,7 @@ do not retarget one onto the other. This stub is the conv contract only. | Math | scalar FMA, fp16 or bf16 in/out, **fp32 acc / mul**. **Not** `fdot2`. **Not** `fdot2.bf16`. gfx1030 LLVM aborts on `llvm.amdgcn.fdot2.bf16.bf16` — BF16 must promote the multiply, not retarget DOT. | | Decode FIR order | extras @ `cafe95ef8` / tip `1ff735`: FIR on **pre-shift** state, then shift (matches fwd). Do not shift-first. | | Fwd out-stride / null-block | extras @ `82b6f183da2a`: write via `stride_o_token`; invalid slot (`null_block_id`) early-returns. | -| State | `state_len = width-1`, typically **3 or 4** for GDN | +| State | `state_len = width-1`, typically **3 or 4** for GDN. Dest @ `b13effcb` gates HIP when spec-decode sizes `state_len != width-1` (Triton fallback). Stay extras. | | Graph | scratch sized by `plan`, **zeroed** for page-commit; **no D2H under capture** | extras decode kernel is cudagraph-safe *as ISA* for **fp16**. Dest GDN