diff --git a/.agents/specs/rocm-fused-norm-rope.md b/.agents/specs/rocm-fused-norm-rope.md new file mode 100644 index 000000000..8bc9a92f7 --- /dev/null +++ b/.agents/specs/rocm-fused-norm-rope.md @@ -0,0 +1,245 @@ +# ROCM-FUSED-NORM-ROPE — `vt::FusedNormRope` on ROCm, so GLM-5.3 reaches a token on `gfx1151` + +Row: `BACKEND-ROCM` +Issue: [#2564](https://github.com/mudler/vllm.cpp/issues/2564) + +## Now + +`ACTIVE`, measured. Base `4fe3852b119e40cda05c0fbcf64d3e2a4796ada2`. + +## Scope + +Register a native `kFusedNormRope` kernel for `DeviceType::kROCM`, and correct +the MLA block's stale comment and refusal message so both causes of the split +A-projection branch are named. + +Out of scope: the other seven missing MLA/DSA ops +(`.agents/specs/rocm-glm53-dsa.md` W1.3). They keep serving from the portable +reference tier and are recorded under `## Owed`. + +## The defect + +`src/vllm/model_executor/layers/attention/mla_attention.cpp:550-551`: + +```cpp +const bool fused_nr = R > 0 && !has_k_rope_norm && MlaFusedNormRopeEnabled() && + vt::OpRegistered(vt::OpId::kFusedNormRope, d.q.device.type); +``` + +`vt::OpRegistered` is a native-only probe by design +(`src/vt/op_provider.cpp:779-803`) and ROCm registers no `kFusedNormRope`. With +every environment variable unset, `fused_nr` is therefore false on `gfx1151`, +the split path row-slices `kv_a_proj_with_mqa`, and a `q8_0` weight has no row +slice — so GLM-5.3's first forward throws. The throw's own comment +(`:628-631`) says the only way to reach it is `VT_MLA_FUSED_NORM_ROPE=0`. That +enumerates the backends that HAVE the op and forgets the ones that do not; the +measured run in #2564 is the counterexample. + +## Why repair 1, and not 2 or 3 + +#2564 prices three repairs. This spec takes the first. + +1. **Port `kFusedNormRope` to ROCm** — taken. Both halves of the composite are + already native ROCm kernels: the latent RMS reduction is + `rocm_rmsnorm.hip`'s `RmsNormRowKernel` and the decoupled-pe rotation is + `rocm_dense_basic.hip`'s `RopeFromCacheK`. The port is the same + composition CUDA already makes, it makes the predicate true honestly, it + runs the work on the device rather than the host, and it decrements the + reference-tier hit count `docs/ROCM.md:60-61` gates on. +2. **Let `fused_nr` consider the reference tier** — rejected. It changes what + "available" means at a shared seam for every backend and every op, and + `src/vt/op_provider.cpp:779-786` states the contract it would break: a + unified accelerator would report every op registered the moment its fallback + installed, and the fused-recipe ladder would stop choosing its portable + composite path. It also keys naturally on host-addressability, and the two + host-addressability predicates on this board answer differently: + `DeviceMemoryIsHostAddressable()` is true (`rocm_backend.hip:371`, which is + what makes the reference tier eligible) while + `HostMemoryIsDeviceAddressable()` is false, because `gfx1151` reports + `pageableMemoryAccess=0` (#2515, measured twice on hardware). A predicate + written against the wrong one of those two reads plausible and answers + backwards on the only board that can test it. +3. **Teach the split path to slice a block-quantized merged row** — rejected + for this wave. It is the largest of the three and it repairs a fallback that + nobody wants taken: the fused arm is bit-identical and one launch cheaper. + It stays owed, because a backend that registers neither op still needs it. + +## Upstream anchors + +Read on the pinned oracle, `~/_git/vllm` @ `5559679229` +(`.agents/upstream-sync.md`). vLLM composes the same two steps, unfused, in the +MLA A-projection: + +- `vllm/model_executor/layers/mla.py:164-165` — `kv_c, k_pe = + kv_lora.split([self.kv_lora_rank, self.qk_rope_head_dim], dim=-1)` and then + `kv_c_normed = self.kv_a_layernorm(kv_c)`. That is the LATENT half of the + fused kernel, and the split it performs is exactly the row split the local + block cannot make on a block-quantized weight. +- `vllm/model_executor/layers/mla.py:175-177` — `self.rotary_emb(positions, + q[..., self.qk_nope_head_dim:], k_pe)`. That is the ROPE half, applied to the + trailing slice of the SAME merged row and to nothing the latent half touches, + which is why fusing the two is arithmetically inert. +- `vllm/model_executor/models/deepseek_v2.py:512-518` — the merged weight is + `ReplicatedLinear(hidden_size, kv_lora_rank + qk_rope_head_dim)`, which fixes + the `[L + R, H]` row shape both arms of the local branch assume. +- `vllm/model_executor/models/deepseek_v2.py:1930` — `class + GlmMoeDsaForCausalLM(DeepseekV2ForCausalLM)`, which + `vllm/model_executor/models/registry.py:117` maps the checkpoint's + architecture string to. **Not `deepseek_v32.py`:** at this pin that file does + not carry the registration, and a reader sent there finds nothing. + +vLLM has no fused kernel for the pair, so upstream is the BEHAVIOURAL reference +and the in-tree CUDA sibling is the STRUCTURAL one, exactly as +`vt::FusedNormRope`'s own contract already records +(`include/vt/ops.h:3388-3410`). + +Ported `file:line`, verbatim composition: + +- `src/vt/cuda/cuda_ops.cu:1163-1245` — `FusedNormRopeKernel`, + `LaunchFusedNormRope`, `FusedNormRopeKernelCuda`. +- `src/vt/rocm/rocm_rmsnorm.hip:67-98` — the `__syncthreads()` shared-memory + tree reduce and the scale loop, reproduced element for element. It is + wavefront-width agnostic, which is the property that makes it portable to a + 64-lane wavefront at all. +- `src/vt/rocm/rocm_dense_basic.hip:607-635` — the cache read and the + neox/gpt-j pairing of `RopeFromCacheK`, reproduced element for element. + +## Design + +One new translation unit, `src/vt/rocm/rocm_mla_fused_norm_rope.hip`, holding +one kernel and its `FusedNormRopeFn` entry point, registered in +`rocm_ops.hip`'s `Registrar` and listed twice in `CMakeLists.txt` (the source +list and the `HIP_ARCHITECTURES` property list). + +Block width stays 256, as `rocm_rmsnorm.hip:41` fixes it and for the reason +stated there: it is four whole wavefronts AND it keeps the reduction order +identical to the CUDA and CPU siblings, which is what keeps the NMSE bar +meaningful. Grid is one block per token, as CUDA's is. + +The two halves address disjoint dims, so the fused output is the composite of +`RmsNorm(x[:, :off])` and `RopeFromCache(x[:, off:])` by construction, not by +inspection. + +## Risks + +- **A wavefront assumption.** The donor reduction uses no warp-level primitive, + so 64-lane wavefronts are safe; the gate below measures it rather than + asserting it. +- **A partial MLA arm hides which half ran.** With the reference tier eligible + a half-ported arm still emits tokens. Mitigated by requiring + `VT_OP_PROVIDER_STATS=1` on every leg and reporting the reference-tier hit + count beside any token (#2505's silent-fallback failure). +- **No speed claim is admissible** from this board for this model while the + hit count is non-zero (`docs/ROCM.md:60-61`). None is made. + +## Tests + +- `tests/vt/test_backend_cross_device.cpp` — a new `FusedNormRope` case, + against the CPU oracle at NMSE <= 5e-4, in both rope styles and both + dtypes, on every backend that registers the op. It is SKIPPED on a build + where no device registers `kFusedNormRope`, which is stated plainly rather + than counted as a pass. +- The e2e leg on `strix:gpu0` is the reachability gate, through the production + entry point (`vllm-cli` -> `vllm_engine_load` -> `ModelRegistry::Forward`), + never a by-hand construction. + +## Gates + +1. `ctest` for the focused unit target on `strix:gpu0`, case AND assertion + counts both read. +2. GLM-5.3 `UD-IQ1_S` through `vllm-cli --device auto` with `VT_CPU_MOE=1`, + greedy, on `strix:gpu0`: generated text printed verbatim, with the + reference-tier hit count beside it. +3. Reachability mutation: remove the `RegisterOp(OpId::kFusedNormRope, + DeviceType::kROCM, ...)` line, REBUILD, rerun the e2e leg. It must throw + #2564's message again. Restore, verify by sha256, rebuild, rerun. + +## Evidence + +All on `strix:gpu0` (`gfx1151`, Radeon 8060S, ROCm `7.2.53211-97f5574fe2`), +under `rc` leases, from tree `b413e323be50822dcbecb30bd61dc90333a416b5`. The +tarball's sha256 was read on both the host and the worker and the two agree +(`cf3841a4983b54b8b972ef06ba43f84034477a60358b984ad6a189ba4bf8210b`). + +**Build.** `ninja rc=0`. `[555/577] Building HIP object +CMakeFiles/vllm.dir/src/vt/rocm/rocm_mla_fused_norm_rope.hip.o` -- the first +compile this TU has ever had. + +**Gate 1, the focused numeric case** (job `6b35b8d3`): `1 test case | 1 passed | +0 failed`, `20 assertions | 20 passed | 0 failed`. + +*The assertion count is the discriminator, not a grep of the trace.* `CAPTURE` +prints only on failure, so grepping the passing run for `ROCM` returns 0 and +proves nothing -- that instrument was useless and is recorded as such. The same +binary's case on a CPU-only build with no accelerator registered runs **2** +assertions. 20 vs 2 is the ROCm arm executing: five `Upload` REQUIREs, two +`CHECK`s and two REQUIREs inside `Nmse`, per rope style. + +**Gate 2, GLM-5.3 e2e through the production entry point** (job `6b35b8d3`): +`VT_CPU_MOE=1 VT_OP_PROVIDER_STATS=1 vllm-cli --model +--device auto --prompt "The capital of France is" --max-tokens 4 +--temperature 0`, `LEG rc=0`, +`prompt_tokens=5 completion_tokens=4 finish_reason=length`, stdout: + +```text + Paris, which is +``` + +`op=114 device=5 selected=vt-native` -- op 114 is `kFusedNormRope`, device 5 is +`kROCM`. The mapping is cross-checked three ways against the run's own named +lines: 29/`ConcatAndCacheMla`, 99/`ConcatMlaNopeRope`, 33/`MlaPrefillAttention`. + +**Five distinct ops served from the portable reference tier** in that run: +`ConcatAndCacheMla`, `ConcatMlaNopeRope`, `MlaPrefillAttention`, +`BatchedMatmul`, `MlaDecodeAttention`. `kFusedNormRope` is not among them. +**No speed number is admissible and none is offered** (`docs/ROCM.md`). + +**Gate 3, the mutation ladder** (job `8b508c69`), each rebuilt and each restored +before the next: + +| Mutation | Build | Result | +|---|---|---| +| M1: delete the `RegisterOp(kFusedNormRope, kROCM)` line | rc=0 | `LEG rc=1`, the #2564 throw reproduced verbatim. KILLED | +| M2: drop the sin term from the rope half | rc=0 | `GATE1 rc=1`, `1 case failed`, `20 assertions | 18 passed | 2 failed`. KILLED | +| restored control | rc=0 | `pre.sha == post.sha` byte-for-byte on both files; `GATE1 rc=0`, 1 case / 20 assertions | + +M1 is simultaneously the RED and the reachability proof. The mutated tree is +behaviourally the base tree at the branch this change repairs, and deleting the +production call site reds the production gate -- which is what +`.agents/reachability.md` asks for and what a by-hand construction cannot show. +Its throw also reads back the corrected message: *"The fused path was not taken +because this backend (rocm) registers NO NATIVE vt::FusedNormRope kernel, and +vt::OpRegistered is a native-only probe that cannot see the portable reference +tier"*. + +**Full cross-device suite:** `29 test cases | 28 passed | 1 failed`, +`80296 assertions | 80295 passed | 1 failed`. The one failure is +`MoeSiluMul matches the CPU oracle`'s bit-exact bf16 `CHECK(got == ref_b)`, +which is the standing red #1954 already tracks on `gfx1200`, now recorded on +`gfx1151` too. This change adds one TU and one registration and touches no MoE +path. It is **not** measured at the base commit on this board, so "pre-existing" +is argued from the absence of code-path overlap rather than from an A/B. + +**Host-side control:** on a CPU-only build of the same tree, +`test_mla_attention_block` runs 21 cases / 2,282,067 assertions green, and the +whole cross-device suite runs 28 cases / 13 assertions green. + +## Stop conditions + +- The board faults or resets in a way that is a property of the board rather + than of this change (#2546 measured 12/12 GPU resets for a gate-sized native + run). Report it as such; do not paper over it. +- A second MLA op turns out to block GENERATION rather than merely make it + slow. Return `NEEDS_DECISION` naming which ops are in which class. + +## Owed + +- The seven remaining MLA/DSA ops on ROCm — `kFusedChain`, `kBatchedMatmul`, + `kConcatAndCacheMla`, `kConcatMlaNopeRope`, `kDsaIndexerLogits`, + `kDsaTopkSelect`, `kGatherMlaCache`, `kMlaDecodeAttention`. Each has a CPU + registration, so each serves from the reference tier on this host-addressable + board and none of them refuses. They are what makes a speed result + inadmissible here. Owned by `BACKEND-ROCM`, recorded in + `.agents/specs/rocm-glm53-dsa.md` W1.5 as a campaign this wave does not open. +- Repair 3 of #2564 — a block-quantized row slice for the split path — for a + backend that registers neither `kFusedNormRope` nor a native alternative. diff --git a/CMakeLists.txt b/CMakeLists.txt index c069396d8..4af073a92 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1721,6 +1721,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gdn_postconv.hip src/vt/rocm/rocm_gdn_scan.hip src/vt/rocm/rocm_gdn_fused.hip + src/vt/rocm/rocm_mla_fused_norm_rope.hip src/vt/rocm/rocm_skinny_gemm.hip src/vt/rocm/rocm_ops.hip) if(VLLM_CPP_HIP_ARCHITECTURES) @@ -1745,6 +1746,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gdn_postconv.hip src/vt/rocm/rocm_gdn_scan.hip src/vt/rocm/rocm_gdn_fused.hip + src/vt/rocm/rocm_mla_fused_norm_rope.hip src/vt/rocm/rocm_skinny_gemm.hip src/vt/rocm/rocm_ops.hip PROPERTIES HIP_ARCHITECTURES "${VLLM_CPP_HIP_ARCHITECTURES}") diff --git a/docs/ROCM.md b/docs/ROCM.md index ac137d578..7c028899a 100644 --- a/docs/ROCM.md +++ b/docs/ROCM.md @@ -70,19 +70,29 @@ Do not use a run with CPU fallbacks as a performance result. | Device enum | [`include/vt/device.h`](../include/vt/device.h) | Compiled and routed through the shared device switch | | Architecture mapping | [`include/vt/rocm/rocm_arch.h`](../include/vt/rocm/rocm_arch.h) | Unit-tested gfx name mapping | | Runtime backend | [`src/vt/rocm/rocm_backend.hip`](../src/vt/rocm/rocm_backend.hip) | Runs on five gfx architectures; managed allocation still needs an integrated-board rerun | -| Operation table | [`src/vt/rocm/rocm_ops.hip`](../src/vt/rocm/rocm_ops.hip) | 44 distinct registered `OpId` values at the recorded count | +| Operation table | [`src/vt/rocm/rocm_ops.hip`](../src/vt/rocm/rocm_ops.hip) | One `Registrar` that names every `OpId` this backend serves natively. Recount it with the command below rather than quoting a number from here | | Kernels | [`src/vt/rocm/`](../src/vt/rocm/) | Dense, GDN, attention, sampling, and the contributor-tested Gemma 4 FP8 MoE path | | Platform | [`src/vllm/platforms/rocm.cpp`](../src/vllm/platforms/rocm.cpp) | Runtime-verified on five gfx architectures | | Attention | [`src/vt/rocm/rocm_paged_attn.hip`](../src/vt/rocm/rocm_paged_attn.hip) | Native paged attention and the SharedK WMMA prefill path | | Build | [`CMakeLists.txt`](../CMakeLists.txt) | `VLLM_CPP_HIP` configuration and build verified on five architectures | | Tests | [`tests/vt/test_rocm_backend.cpp`](../tests/vt/test_rocm_backend.cpp) | Runtime cases pass; managed-allocation cases remain pending | -Recount registered operations before you quote the total: +Recount registered operations before you quote the total. The scan must not +depend on where the argument list wraps: several calls in `rocm_ops.hip` break +the line after `RegisterOp(`, and a line-based `grep` never sees `RegisterOp(` +and `OpId::` together on those (#2573). Read the whole file and match across +newlines: ```sh -grep -rho 'RegisterOp(OpId::[A-Za-z0-9_]*' src/vt// | sort -u | wc -l +grep -rhoz 'RegisterOp([[:space:]]*OpId::[A-Za-z0-9_]*[[:space:]]*,[[:space:]]*DeviceType::kROCM' \ + src/vt/rocm/ | tr '\0' '\n' | grep -o 'OpId::[A-Za-z0-9_]*' | sort -u | wc -l ``` +Substitute the `DeviceType::` value for another backend. Naming the device in the +pattern is what keeps the count answering the question a reader asked -- the +previous command counted every `RegisterOp(OpId::` line in the directory +regardless of which device it registered for. + ## Hardware notes | Hardware | Architecture | Memory | Current path | diff --git a/docs/USAGE.md b/docs/USAGE.md index b0c646552..56898ee5f 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -655,7 +655,7 @@ repository in this project's history. | GLM-5.3-Flash FP8 source | `model-000{01..62}-of-00062.safetensors` | 328,326,771,576 bytes total (305.78 GiB) | `zai-org/GLM-5.3-Flash` @ `main`, read 2026-08-26 | Owed: no byte of payload has been fetched, so no local hash exists to state, and an unauthenticated tree hash is not a pin here | Declared source of `scripts/convert-glm5-next-gguf.py`. Only the safetensors HEADERS were read, by HTTP RANGE over all 62 shards: 76,108 tensors, `F8_E4M3` block-quantized at `weight_block_size: [128, 128]` with `weight_scale_inv` companions, plus BF16 and F32 scales | **Nothing has been converted.** The download needs explicit developer authority and a box with room for 305.78 GiB of source and ~100.35 GiB of output at once; owed as O7 on [#2011](https://github.com/mudler/vllm.cpp/issues/2011). The revision is a branch name and not a commit, which is NOT a pin: it is what was read, and W7b re-reads and records the commit when it stages the bytes | | GLM-5.3-Flash GGUF | `GLM-5.3-Flash-UD-Q2_K_XL-0000{1..4}-of-00004.gguf` | 108,720,071,427 bytes total (101.2535 GiB) across four shards; 1412 tensors | `unsloth/GLM-5.3-Flash-GGUF` @ `d425e572fb9686125831f476129e51cea34bc5b4`, path `UD-Q2_K_XL`, staged 2026-08-28 | Owed for this row: the shards are staged and were sha256-verified when they were fetched, but **W5c consumed only the four GGUF HEADERS** and states no hash of its own. W7b ([#2225](https://github.com/mudler/vllm.cpp/issues/2225)) records the per-shard sha256 alongside the load it measures | **LOADS on `--device cpu`, and the engine's multi-KV guard no longer refuses above the model's forward** ([#2348](https://github.com/mudler/vllm.cpp/issues/2348)). **A MATERIALIZED LOAD EXISTS** -- driven at this artifact on `dgx:gpu0` 2026-08-30, all four shards load and the engine sizes its caches in under 26 minutes wall ([#2343](https://github.com/mudler/vllm.cpp/issues/2343)). At that change the first step threw at the `multi_kv` guard above the model's own hook; W5b-2c ([#2348](https://github.com/mudler/vllm.cpp/issues/2348)) writes the consuming forward that guard was waiting for and it no longer fires for this model. **THIS ARTIFACT GENERATES COHERENT TEXT, and peak RSS is MEASURED** as of [#2241](https://github.com/mudler/vllm.cpp/issues/2241). On `dgx:gpu0` 2026-08-30, in the SHIPPED configuration with no diagnostic env set, `vllm-cli --device cpu --max-tokens 2` at the prompt `The capital of France is` emits ` Paris.` at `rc=0`, and `VmHWM` peaks at 104,792,300 kB = 99.94 GiB. Two instrumented `thor:gpu0` runs the same day supply the bisect: four tokens read ` Paris. Paris is`, the prefill top-5 is ` Paris` (16.427), ` one`, ` located`, ` known`, ` a` at margin 1.279, and none of 180 per-layer readings over four steps carries a NaN. The first attempt emitted token id 0 eight times, because the loader repacked this file's 346 q8_0 tensors into the i8mm interleave that the host bridge reads as plain blocks (spec `## Owed` O30). **No speed number is claimed, and the earlier ones are void** -- they were taken from an all-NaN forward. The GB10 arm is the one measured above. The GGUF arm of `load_weights` resolves all 1383 backbone tensors of this file (W5c, [#2242](https://github.com/mudler/vllm.cpp/issues/2242)); `blk.45`, the multi-token-prediction block, is read, counted and DROPPED, as the transformers reference does. `ModelRegistry::Forward` dispatches to the model as of W5b-2b ([#2337](https://github.com/mudler/vllm.cpp/issues/2337)), which bridges ONE decoder layer at a time out of the block-resident tower and decodes only the 8 of 288 experts a token selects — a float tower is 426.72 GiB against ~119.63 GiB usable. **A MATERIALIZED LOAD NOW EXISTS**: driven at this artifact on `dgx:gpu0` 2026-08-30, all four shards load and the engine sizes its caches in under 26 minutes wall. **NO TOKEN WAS GENERATED** — the first step throws at the `multi_kv` guard above the model's own hook ([#2343](https://github.com/mudler/vllm.cpp/issues/2343), [#2068](https://github.com/mudler/vllm.cpp/issues/2068)) — and **peak RSS and speed are still unmeasured**, because the staging run did not sample them. The vision tower (a separate `mmproj-BF16.gguf`) and the safetensors arm still refuse by name, as does a multi-request step; a non-CPU queue is admitted as of W9c-3a ([#2464](https://github.com/mudler/vllm.cpp/issues/2464)) for the routed-expert GEMM alone, and a device that is neither CPU nor CUDA is refused by name; **the KV-cache spec does not**, as of W5 ([#2223](https://github.com/mudler/vllm.cpp/issues/2223)), which publishes its three groups through the production factory hook | **The earlier row here said `none exists`, and that was true when it was written (2026-08-26) and is not now.** "UD-Q2_K_XL" names a TARGET AVERAGE and not a format: the census over all 1412 tensors is F32 638, Q8_0 346, Q5_K 181, Q6_K 117, IQ2_XS 82, IQ3_XXS 41, IQ4_XS 3, Q2_K 2, Q4_K 1, Q3_K 1 — **two** Q2_K tensors in a file named Q2_K. It fits `dgx:gpu0` only because IQ2_XS and IQ4_XS keep their blocks ([#2247](https://github.com/mudler/vllm.cpp/issues/2247)); both now have a CUDA keep-quant kernel too ([#2260](https://github.com/mudler/vllm.cpp/issues/2260)), so the expert GEMM no longer drains the stream to the host and the fused seam no longer throws. W9c-3a ([#2464](https://github.com/mudler/vllm.cpp/issues/2464)) then built a device arm for this artifact's routed-expert GEMM and MEASURED it end to end, where it **SEGFAULTED**: both `--device cuda` legs on `dgx:gpu0` died with rc=139 emitting no token, reproducibly (spec O46). The split is therefore OPT-IN and defaults OFF, so `--device cuda` refuses exactly as it did before. **Use `--device cpu`** -- measured on that artifact it emits ` Paris.` at rc=0, 1176 s wall of which 169 s is generation. Every OTHER primitive of this model is still a host reference on an interposed CPU queue (spec O43), so what `--device cuda` reaches is one arm of eleven and not a device arm. **A materialized load NOW exists and a token still does not** — `dgx:gpu0` 2026-08-30 ([#2343](https://github.com/mudler/vllm.cpp/issues/2343)): all four shards load and the engine sizes its caches, then the first step throws at the `multi_kv` guard above the model's own hook. **Peak RSS and speed remain unmeasured** | | GLM-5.3-Flash config | `config.json` | 69,416 bytes | `zai-org/GLM-5.3-Flash` @ `main`, read 2026-08-27 | sha256 `bb8f01c42cb92a52ca72e65afb4d5bd8d11aef083cd210e8de25dfb904f23e9f` | The ONLY byte of this checkpoint any change on this row has consumed. Checked in verbatim as `tests/vllm/models/fixtures/glm5_next/config.json` and used as W1's gate fixture, so the config layer is gated against what the checkpoint says rather than against what a port's author believed it says | **Arms refused by name:** the SAFETENSORS one, which is what this row is, because every published safetensors artifact of this model exceeds every device this project owns. `Glm5NextForConditionalGeneration` is REGISTERED, its config RESOLVES, and the GGUF arm both loads and forwards ([#2067](https://github.com/mudler/vllm.cpp/issues/2067), [#2242](https://github.com/mudler/vllm.cpp/issues/2242), [#2337](https://github.com/mudler/vllm.cpp/issues/2337)). The revision is a branch name and not a commit, which is NOT a pin for the WEIGHTS; for this one file the sha256 above is the pin | -| GLM-5.3 GGUF (`glm-dsa`) all six shards | `GLM-5.3-UD-IQ1_S-0000{1..6}-of-00006.gguf` | 9,428,677 B (shard 1, metadata only, 0 tensors) and 49,968,868,928 B (shard 2); the six shards total 216,715,365,893 B = 201.83 GiB across 1809 tensors | `unsloth/GLM-5.3-GGUF` @ `346b3591c7f28d1a23716f97a065ecf12ec14771`, path `UD-IQ1_S`, staged 2026-08-30, completed and verified 2026-08-31 | shard 1 `ff3adab0853dfb00bdf3889ec3f5556196f56b65783115720d57767bbd760dd9`; shard 2 `659d04cf4fc0b6026944f34c0b590a635803bff06c1775361e28490db7b168f8`; shard 3 `433302bac0e2d54da64c7c2f28509fa1b235aeccdf5b215a8a446ebaad1b5b27`; shard 4 `d0a6f19452d5b5cd498e1eb8fbe856e00aed7da1f80c27c095301eabe81e9bc1`; shard 5 `2ea1537ffab40fa8b8584a8647ec10fbaa6199dfed45e4019b822da2b319db37`; shard 6 `42a76ef04ffc5e321e1240f4e572b6fa6fc3315da5bea22fb598d7460db210fe`. **All six are complete and each was hashed TWICE** — once by the fetch script as it landed and once independently off the same share afterwards — and the two readings agree. **The DERIVED metadata shard has a hash of its own:** `scripts/glm-dsa-write-indexer-types.py` run against the staged shard 1 with `zai-org/GLM-5.3`'s own `config.json` produces a 9,428,810-byte file, 64 keys becoming 65, 21 `full` of 78, sha256 `b3e9838651a5c279533c98390ab4bc03cf1d8c176d5be0754180f07d9ed85c01`, reproduced identically by three independent runs. **That is a DERIVED artifact and must never be quoted as `unsloth/GLM-5.3-GGUF`'s shard 1** | **THIS ARTIFACT GENERATES THROUGH THE EXPERT-STREAMING LANE: `The capital of France is` -> ` Paris`.** On `dgx:gpu0` (GB10, 20 cores, 119 GB, compute capability 12.1) under an `rc` lease, 2026-08-31, a build with `-DVLLM_CPP_CUDA_ARCHITECTURES=121a -DVLLM_CPP_FLASH_ATTN=ON` and CUTLASS 4.5.0: `VT_MOE_EXPERT_STREAM=1 VT_MOE_EXPERT_STREAM_SLOTS=4096 vllm-cli --model --device cuda --prompt "The capital of France is" --max-tokens 1 --temperature 0` returns `rc=0`, `prompt_tokens=5 completion_tokens=1`, and seven bytes of stdout: a space, `Paris`, a newline. Wall 1154 s for the process, `generate` 852.330 s, `VmHWM` 60,512,268 kB = 57.71 GiB against 119.631 GiB of device and 201.83 GiB of artifact. **The lane's own counters are the streaming evidence, and they include one number that must travel with them:** `[expert-stream] ON slots=4096 slot_bytes=6684672 resident=25.50 GiB`, then `steps=1 hits=0 misses=6399 evictions=0 fills=4096 bytes=13939408896 exhausted=2303 advised=0`. 4096 slices were paged out of the file into slots and 12.98 GiB moved through them with zero evictions, and the 187.312 GiB of towers were never materialized — but the step needed 6399 distinct slices, so **2303 of them (36%) were read in place out of the mapping instead of streamed**. That is a PREFILL working set exceeding any slot budget by construction (spec R2, O34), it is counted rather than silent, and no figure here may be quoted as a fully-streamed step. **No speed number is claimed:** one token, a CIFS-backed artifact, and 2303 in-place fallbacks in the measurement. `--device cpu` on the same box and artifact also emits ` Paris` (`rc=0`, `generate` 950.249 s, `VmHWM` 44.46 GiB), and **that arm does NOT stream at all** — a CPU queue builds no slot lane, so every routed-expert slice is read in place. On `thor:gpu0` (sm_110a) the CUDA arm cannot reach a token: MLA prefill on this family IS FlashAttention, the vendored FA2 covers `8.0,8.6,8.7,8.9,12.0a,12.1a`, and sm_110a is outside it. Also gated on a complete synthetic model of the same shape: `test_glm_moe_dsa_gguf_load.cpp` 5 cases / 228 assertions, `test_glm_moe_dsa_forward.cpp` 7 / 5258, `test_glm_moe_dsa_schedule.cpp` 12 / 533, and the real file's census from its headers (`test_glm_moe_dsa_gguf_census.cpp` 3 / 3831): 1809 tensors, 228 expert towers at 187.312 GiB, 1581 resident at 14.511 GiB, largest per-expert slice 6,684,672 B. What the forward still refuses BY NAME is a step in which any request RESUMES while its selection PRUNES — that needs the indexer KV side cache `KV-DSV4-MULTICACHE` owns (spec O4, #1925/#2323), so a FIRST token on a fresh prompt is reachable and a SECOND is not — and sparse prefill (spec O6) is still W6's. No speed axis has a denominator (spec O10) | **THIS FILE CANNOT BE FED AS PUBLISHED**, and that is a property of the file rather than of the port: its 64 metadata keys carry neither `glm-dsa.attention.indexer.types` nor `index_topk_freq`/`index_skip_topk_offset`, so it states its per-layer indexer schedule nowhere, and it broadcasts `indexer.*` onto all 79 blocks while the checkpoint ships them on 22. The loader refuses it by name rather than substituting llama.cpp's hardcoded table (spec D3). **The repair is one command and it rewrites the 9.4 MB metadata shard only:** `scripts/glm-dsa-write-indexer-types.py --shard --from-config --out /GLM-5.3-UD-IQ1_S-00001-of-00006.gguf`, with the five payload shards hard-linked beside the output, then `--model` that directory's shard 1. It transcribes the schedule from the model author's own `config.json` and derives nothing; the result is a DERIVED artifact with its own sha256 and is not `unsloth/GLM-5.3-GGUF`. **Build requirements this model does not degrade past:** `--device cuda` (the expert-streaming lane is not built on a CPU queue, and the towers would then be read in place out of a 201.83 GiB mmap), and a build with the vendored FlashAttention-2, which needs CUTLASS headers and an arch in `8.0,8.6,8.7,8.9,12.0a,12.1a` — MLA prefill IS FlashAttention here and has no fallback below it. **On ROCm `gfx1151` (`strix:gpu0`, Radeon 8060S) this artifact LOADS and DOES NOT GENERATE, as of [#2562](https://github.com/mudler/vllm.cpp/pull/2562).** The route is `VT_CPU_MOE=1 vllm-cli --model --device auto` -- `auto` because no `--device` value names ROCm ([#2505](https://github.com/mudler/vllm.cpp/issues/2505)), and `cpu_moe` because `--fit`'s default placement leaves 22 layers on a device whose keep-quant set cannot hold their IQ1_S towers ([#2565](https://github.com/mudler/vllm.cpp/issues/2565)). All 1809 tensors resolve, all 228 routed-expert towers stay compressed, 11.620 GiB is paged in (at 11.5 MiB/s off the CIFS share the artifact lives on, which is a property of the share), and the engine auto-fits `max_model_len` to 8192 against 256 blocks of 32 tokens. **NO TOKEN COMES OUT:** the first forward throws in the MLA block, because `vt::OpRegistered(kFusedNormRope, ...)` is false on ROCm and cannot see the reference tier, so the split A-projection path is taken and refuses this checkpoint's block-quantized `kv_a_proj_with_mqa` ([#2564](https://github.com/mudler/vllm.cpp/issues/2564)). The streamed-expert lane is NOT what serves the towers here and cannot be: `pageableMemoryAccess` is 0 on this board, so `host_memory_is_device_addressable()` is false ([#2515](https://github.com/mudler/vllm.cpp/issues/2515)). No speed number is admissible from this board for this model. **Arms refused by name:** the SAFETENSORS one, permanently (spec D1 — 703.74 GiB across 141 shards, no streaming loader, no MoE block-fp8 rung), and `UD-IQ1_M`, which refuses at file open because `IQ1_M` (ggml id 29) has no reader traits (spec O3) | +| GLM-5.3 GGUF (`glm-dsa`) all six shards | `GLM-5.3-UD-IQ1_S-0000{1..6}-of-00006.gguf` | 9,428,677 B (shard 1, metadata only, 0 tensors) and 49,968,868,928 B (shard 2); the six shards total 216,715,365,893 B = 201.83 GiB across 1809 tensors | `unsloth/GLM-5.3-GGUF` @ `346b3591c7f28d1a23716f97a065ecf12ec14771`, path `UD-IQ1_S`, staged 2026-08-30, completed and verified 2026-08-31 | shard 1 `ff3adab0853dfb00bdf3889ec3f5556196f56b65783115720d57767bbd760dd9`; shard 2 `659d04cf4fc0b6026944f34c0b590a635803bff06c1775361e28490db7b168f8`; shard 3 `433302bac0e2d54da64c7c2f28509fa1b235aeccdf5b215a8a446ebaad1b5b27`; shard 4 `d0a6f19452d5b5cd498e1eb8fbe856e00aed7da1f80c27c095301eabe81e9bc1`; shard 5 `2ea1537ffab40fa8b8584a8647ec10fbaa6199dfed45e4019b822da2b319db37`; shard 6 `42a76ef04ffc5e321e1240f4e572b6fa6fc3315da5bea22fb598d7460db210fe`. **All six are complete and each was hashed TWICE** — once by the fetch script as it landed and once independently off the same share afterwards — and the two readings agree. **The DERIVED metadata shard has a hash of its own:** `scripts/glm-dsa-write-indexer-types.py` run against the staged shard 1 with `zai-org/GLM-5.3`'s own `config.json` produces a 9,428,810-byte file, 64 keys becoming 65, 21 `full` of 78, sha256 `b3e9838651a5c279533c98390ab4bc03cf1d8c176d5be0754180f07d9ed85c01`, reproduced identically by three independent runs. **That is a DERIVED artifact and must never be quoted as `unsloth/GLM-5.3-GGUF`'s shard 1** | **THIS ARTIFACT GENERATES THROUGH THE EXPERT-STREAMING LANE: `The capital of France is` -> ` Paris`.** On `dgx:gpu0` (GB10, 20 cores, 119 GB, compute capability 12.1) under an `rc` lease, 2026-08-31, a build with `-DVLLM_CPP_CUDA_ARCHITECTURES=121a -DVLLM_CPP_FLASH_ATTN=ON` and CUTLASS 4.5.0: `VT_MOE_EXPERT_STREAM=1 VT_MOE_EXPERT_STREAM_SLOTS=4096 vllm-cli --model --device cuda --prompt "The capital of France is" --max-tokens 1 --temperature 0` returns `rc=0`, `prompt_tokens=5 completion_tokens=1`, and seven bytes of stdout: a space, `Paris`, a newline. Wall 1154 s for the process, `generate` 852.330 s, `VmHWM` 60,512,268 kB = 57.71 GiB against 119.631 GiB of device and 201.83 GiB of artifact. **The lane's own counters are the streaming evidence, and they include one number that must travel with them:** `[expert-stream] ON slots=4096 slot_bytes=6684672 resident=25.50 GiB`, then `steps=1 hits=0 misses=6399 evictions=0 fills=4096 bytes=13939408896 exhausted=2303 advised=0`. 4096 slices were paged out of the file into slots and 12.98 GiB moved through them with zero evictions, and the 187.312 GiB of towers were never materialized — but the step needed 6399 distinct slices, so **2303 of them (36%) were read in place out of the mapping instead of streamed**. That is a PREFILL working set exceeding any slot budget by construction (spec R2, O34), it is counted rather than silent, and no figure here may be quoted as a fully-streamed step. **No speed number is claimed:** one token, a CIFS-backed artifact, and 2303 in-place fallbacks in the measurement. `--device cpu` on the same box and artifact also emits ` Paris` (`rc=0`, `generate` 950.249 s, `VmHWM` 44.46 GiB), and **that arm does NOT stream at all** — a CPU queue builds no slot lane, so every routed-expert slice is read in place. On `thor:gpu0` (sm_110a) the CUDA arm cannot reach a token: MLA prefill on this family IS FlashAttention, the vendored FA2 covers `8.0,8.6,8.7,8.9,12.0a,12.1a`, and sm_110a is outside it. Also gated on a complete synthetic model of the same shape: `test_glm_moe_dsa_gguf_load.cpp` 5 cases / 228 assertions, `test_glm_moe_dsa_forward.cpp` 7 / 5258, `test_glm_moe_dsa_schedule.cpp` 12 / 533, and the real file's census from its headers (`test_glm_moe_dsa_gguf_census.cpp` 3 / 3831): 1809 tensors, 228 expert towers at 187.312 GiB, 1581 resident at 14.511 GiB, largest per-expert slice 6,684,672 B. What the forward still refuses BY NAME is a step in which any request RESUMES while its selection PRUNES — that needs the indexer KV side cache `KV-DSV4-MULTICACHE` owns (spec O4, #1925/#2323), so a FIRST token on a fresh prompt is reachable and a SECOND is not — and sparse prefill (spec O6) is still W6's. No speed axis has a denominator (spec O10) | **THIS FILE CANNOT BE FED AS PUBLISHED**, and that is a property of the file rather than of the port: its 64 metadata keys carry neither `glm-dsa.attention.indexer.types` nor `index_topk_freq`/`index_skip_topk_offset`, so it states its per-layer indexer schedule nowhere, and it broadcasts `indexer.*` onto all 79 blocks while the checkpoint ships them on 22. The loader refuses it by name rather than substituting llama.cpp's hardcoded table (spec D3). **The repair is one command and it rewrites the 9.4 MB metadata shard only:** `scripts/glm-dsa-write-indexer-types.py --shard --from-config --out /GLM-5.3-UD-IQ1_S-00001-of-00006.gguf`, with the five payload shards hard-linked beside the output, then `--model` that directory's shard 1. It transcribes the schedule from the model author's own `config.json` and derives nothing; the result is a DERIVED artifact with its own sha256 and is not `unsloth/GLM-5.3-GGUF`. **Build requirements this model does not degrade past:** `--device cuda` (the expert-streaming lane is not built on a CPU queue, and the towers would then be read in place out of a 201.83 GiB mmap), and a build with the vendored FlashAttention-2, which needs CUTLASS headers and an arch in `8.0,8.6,8.7,8.9,12.0a,12.1a` — MLA prefill IS FlashAttention here and has no fallback below it. **On ROCm `gfx1151` (`strix:gpu0`, Radeon 8060S) this artifact NOW GENERATES TEXT, as of [#2572](https://github.com/mudler/vllm.cpp/pull/2572): `The capital of France is` -> ` Paris, which is`.** It loaded but emitted nothing between [#2562](https://github.com/mudler/vllm.cpp/pull/2562) and that change. The route is `VT_CPU_MOE=1 vllm-cli --model --device auto` -- `auto` because no `--device` value names ROCm ([#2505](https://github.com/mudler/vllm.cpp/issues/2505)), and `cpu_moe` because `--fit`'s default placement leaves 22 layers on a device whose keep-quant set cannot hold their IQ1_S towers ([#2565](https://github.com/mudler/vllm.cpp/issues/2565)). All 1809 tensors resolve, all 228 routed-expert towers stay compressed, 11.620 GiB is paged in (at 11.5 MiB/s off the CIFS share the artifact lives on, which is a property of the share), and the engine auto-fits `max_model_len` to 8192 against 256 blocks of 32 tokens. **TOKENS NOW COME OUT, and what serves them must travel with them.** The first forward used to throw in the MLA block, because `vt::OpRegistered(kFusedNormRope, ...)` was false on ROCm and cannot see the reference tier, so the split A-projection path was taken and refused this checkpoint's block-quantized `kv_a_proj_with_mqa` ([#2564](https://github.com/mudler/vllm.cpp/issues/2564)). Registering a native ROCm `kFusedNormRope` makes that predicate true, and the run completes: measured 2026-09-02, `rc` job `6b35b8d3-be7f-4d71-abf4-0f0bd72bb643`, `VT_CPU_MOE=1 VT_OP_PROVIDER_STATS=1 vllm-cli --model --device auto --prompt "The capital of France is" --max-tokens 4 --temperature 0` returns `rc=0`, `prompt_tokens=5 completion_tokens=4 finish_reason=length`, and prints ` Paris, which is`. **FIVE ops ran on the portable CPU reference tier in that run** -- `ConcatAndCacheMla`, `ConcatMlaNopeRope`, `MlaPrefillAttention`, `BatchedMatmul` and `MlaDecodeAttention` -- and the `kFusedNormRope` this change adds is NOT one of them (`op=114 device=5 selected=vt-native`). **NO SPEED NUMBER IS ADMISSIBLE from this run and none is offered:** `docs/ROCM.md` disqualifies any performance result with a non-zero reference-tier hit count, and this run has five. The 3516.719 s the harness printed for four tokens is recorded here only as the cost of a host-tier MLA arm, never as a throughput result. The streamed-expert lane is NOT what serves the towers here and cannot be: `pageableMemoryAccess` is 0 on this board, so `host_memory_is_device_addressable()` is false ([#2515](https://github.com/mudler/vllm.cpp/issues/2515)). No speed number is admissible from this board for this model. **Arms refused by name:** the SAFETENSORS one, permanently (spec D1 — 703.74 GiB across 141 shards, no streaming loader, no MoE block-fp8 rung), and `UD-IQ1_M`, which refuses at file open because `IQ1_M` (ggml id 29) has no reader traits (spec O3) | | GLM-5.3 config | `config.json` | 29,464 bytes | `zai-org/GLM-5.3` @ `935644c05e76fc198714f4cca449fd8b970ff6d7` | Committed verbatim in-tree as `tests/vllm/models/glm_moe_dsa_config_glm53.inc`, so the config layer is gated against what the checkpoint says rather than against what a port's author believed it says | It is the ONLY authoritative source of the 78-entry `indexer_types` list — 21 `full`, at layers {0,1,2} and every fourth from 6 to 74 — which three independent derivations agree on bit for bit (the list itself, vLLM's rule at `deepseek_v2.py:1097-1101`, and llama.cpp's `GLM_5_2_DEFAULT_INDEXER_TYPES`) | The GGUF above does not carry this list, which is why it cannot be fed as published | | Qwen3.5-0.8B (Tenstorrent P150 arm) | `model.safetensors-00001-of-00001.safetensors` | 1,746,942,600 bytes | `Qwen/Qwen3.5-0.8B` @ `2fc06364715b967f1860aea9cf38778875588b17`, authorized 2026-08-23 | `04b1c301231dd422b8860db31311ab2721511346a32cb1e079c4c4e5f1fe4696` (non-quantized; hashed anyway from the local bytes the gates and the eager profile consumed) | bf16 on the Tenstorrent P150: the sacred greedy pair, both ambient legs, and the #1715/#2107 profile legs all ran from this snapshot | **Arms refused by name:** GGUF k-quant arms on TT — no TT kernels exist for them, refused at load; Qwen3.8-27B on TT — no arm fits the P150 (bf16 53.8 GB), refused at load | | dots3-note bf16 language tower | `model-000{01..131}-of-00131.safetensors` | 561,371,869,568 bytes total (522.82 GiB), of which the MoE is 545,823,175,680 | `dots-studio/dots3-note-prev` @ `1e1e7b0cd37a3a48a6c8d7fa55d5f9d14377006b` | Owed: **no tensor byte has been fetched**, so no local hash exists to state, and an unauthenticated tree hash is not a pin here | The bf16 text tower this port loads: 46 backbone layers, both MLA geometries, and since W5 the 45 MoE layers — the ungrouped noaux_tc router at 256/8 plus one shared expert at `moe_intermediate_size * n_shared_experts` = 1536. Everything except `mlp.gate.e_score_correction_bias` is BF16; that one is F32, on both sides | **Nothing has ever loaded these bytes.** The tower alone is 522.82 GiB against a 122 GiB ceiling on the largest host this project reaches (spec §6.2), so the arm is representable and unfeedable, and the e2e gate is an OPEN GAP by construction. GGUF k-quants are refused by name (W9). The 19-tensor nextn tail is a NAMED W10 deferral rather than a refusal since #2176 | diff --git a/src/vllm/model_executor/layers/attention/mla_attention.cpp b/src/vllm/model_executor/layers/attention/mla_attention.cpp index c05003120..ac7c9a007 100644 --- a/src/vllm/model_executor/layers/attention/mla_attention.cpp +++ b/src/vllm/model_executor/layers/attention/mla_attention.cpp @@ -626,20 +626,39 @@ void ForwardMlaAttentionBlock(Dev d, const MlaBlockDims& dims, const MlaBlockWei // HAS NO ROW SLICE. `Tensor::Slice` offsets by `stride[dim] * // SizeOf(dtype)` and `SizeOf` refuses a block dtype outright, so this // would throw four frames deeper with a message about dtypes rather than - // about the operator's switch. `vt::FusedNormRope` is registered on CPU - // (`cpu_ops.cpp`) and CUDA (`cuda_ops.cu`) and is default-ON, so the only - // way here on a keep-quant checkpoint is `VT_MLA_FUSED_NORM_ROPE=0`. + // about the operator's switch. + // + // TWO things can put a keep-quant checkpoint here, and the message says + // WHICH by reading the predicate terms back rather than naming one and + // hoping (#2564). The previous text named only the environment override, + // because it enumerated the backends that HAVE `vt::FusedNormRope` — CPU + // (`cpu_ops.cpp`) and CUDA (`cuda_ops.cu`) — and forgot the ones that do + // not. `vt::OpRegistered` is a NATIVE-ONLY probe by design + // (`src/vt/op_provider.cpp:779-803`), so a backend with no native kernel + // takes this branch with every environment variable unset even where the + // portable reference tier would have served the op. That is what #2564 + // measured on `gfx1151` with GLM-5.3, and a reader sent looking for an + // unset variable finds nothing. if (vt::IsBlockQuant(w_kva.dtype)) { + const bool env_off = !MlaFusedNormRopeEnabled(); + std::string why = + env_off ? "VT_MLA_FUSED_NORM_ROPE=0 is set, which disables the fused op" + : (std::string("this backend (") + vt::DeviceTypeName(d.q.device.type) + + ") registers NO NATIVE vt::FusedNormRope kernel, and " + "vt::OpRegistered is a native-only probe that cannot see " + "the portable reference tier"); throw std::invalid_argument( "MLA block: the split A-projection path needs vt::FusedNormRope to " "read the merged [kv_lora_rank + qk_rope_head_dim] row, because a " "BLOCK-QUANTIZED kv_a_proj_with_mqa (" + std::string(vt::Name(w_kva.dtype)) + ") has no row slice — a quant block spans whole rows and " - "vt::SizeOf refuses a per-element size for it. This is reachable " - "only with VT_MLA_FUSED_NORM_ROPE=0 on a keep-quant MLA " - "checkpoint; unset it, or load this model with an expanded " - "residency"); + "vt::SizeOf refuses a per-element size for it. The fused path was " + "not taken because " + + why + + (env_off ? "; unset it, or load this model with an expanded residency" + : "; port kFusedNormRope to this backend, or load this " + "model with an expanded residency")); } Tensor kv_c_t = kv_c.t(), k_pe_t = k_pe.t(); vt::MatmulBT(d.q, kv_c_t, hidden, w_kva.Slice(0, 0, L)); diff --git a/src/vt/rocm/rocm_mla_fused_norm_rope.hip b/src/vt/rocm/rocm_mla_fused_norm_rope.hip new file mode 100644 index 000000000..528334d76 --- /dev/null +++ b/src/vt/rocm/rocm_mla_fused_norm_rope.hip @@ -0,0 +1,181 @@ +// ROCm `vt::FusedNormRope` (kFusedNormRope) — the Tier-A2+A5 MLA norm-rope fold +// (ROCM-FUSED-NORM-ROPE, issue #2564). +// +// WHY THIS OP IS NOT OPTIONAL ON THIS BACKEND. Every other missing MLA op falls +// through to the portable reference tier and merely runs on the host. This one +// does not, because `src/vllm/model_executor/layers/attention/mla_attention.cpp` +// BRANCHES on `vt::OpRegistered(kFusedNormRope, device)` before it ever calls +// the op, and `OpRegistered` is a native-only probe by design +// (`src/vt/op_provider.cpp:779-803`). A backend that does not register this +// kernel therefore takes the SPLIT A-projection path, which row-slices +// `kv_a_proj_with_mqa` — and a block-quantized weight has no row slice. So the +// absence of this kernel is a REFUSAL on a keep-quant MLA checkpoint, not a +// slowdown. That is what #2564 measured on `strix:gpu0` with GLM-5.3. +// +// UPSTREAM (pinned oracle `~/_git/vllm` @ `5559679229`). vLLM has no fused +// kernel for this pair; it composes the two steps, and this kernel is that +// composition: +// vllm/model_executor/layers/mla.py:164-165 — `kv_c, k_pe = kv_lora.split( +// [kv_lora_rank, qk_rope_head_dim], dim=-1)` then `self.kv_a_layernorm(kv_c)` +// vllm/model_executor/layers/mla.py:175-177 — `self.rotary_emb(positions, ..., +// k_pe)` over the trailing slice of the SAME merged row +// vllm/model_executor/models/deepseek_v2.py:512-518 — the merged weight is +// [kv_lora_rank + qk_rope_head_dim, hidden_size], which fixes the row shape +// The two halves address DISJOINT dims, so fusing them is arithmetically inert. +// +// PORTED FROM: src/vt/cuda/cuda_ops.cu:1163-1245 (`FusedNormRopeKernel`, +// `LaunchFusedNormRope`, `FusedNormRopeKernelCuda`), which is itself the +// composition of two kernels this backend ALREADY has natively: +// * the latent RMS reduce + scale — src/vt/rocm/rocm_rmsnorm.hip:67-98 +// (`RmsNormRowKernel`), reproduced element for element. It is a +// __syncthreads() tree over shared memory with NO warp-level primitive, so +// it is wavefront-width agnostic; that is the property that makes it +// portable to AMD's 64-lane wavefront unchanged, and it is the reason +// `kBlock` stays 256 rather than being raised to a wavefront multiple — +// 256 is already four whole wavefronts AND it keeps the reduction ORDER +// identical to the CUDA and CPU siblings, which is what makes the NMSE bar +// mean anything. +// * the cache read + neox/gpt-j rotation — src/vt/rocm/rocm_dense_basic.hip: +// 607-635 (`RopeFromCacheK`), reproduced element for element, specialized +// to the single-vector rank-1-positions case the MLA decoupled-pe slice is. +// +// Gate: the FusedNormRope arm of tests/vt/test_backend_cross_device.cpp (NMSE +// <= 5e-4 against the CPU oracle, both rope styles, f32 and bf16), and the +// GLM-5.3 e2e leg on `strix:gpu0` through `vllm-cli`. +#include +#include + +#include +#include +#include + +#include "vt/ops.h" +#include "vt/rocm/rocm_device_bind.h" + +namespace vt::rocm { +namespace { + +// cuda_ops.cu:24 / rocm_rmsnorm.hip:41. See the file header for why it is not +// widened to a wavefront multiple. +constexpr int kBlock = 256; + +inline void Check(hipError_t err, const char* what) { + if (err != hipSuccess) { + throw std::runtime_error(std::string("vt rocm fused_norm_rope: ") + what + ": " + + hipGetErrorString(err)); + } +} +inline hipStream_t AsStream(const Queue& q) { return static_cast(q.handle); } + +__device__ inline float Load(const float* p, int64_t i) { return p[i]; } +__device__ inline float Load(const __hip_bfloat16* p, int64_t i) { + return __bfloat162float(p[i]); +} +__device__ inline void Store(float* p, int64_t i, float v) { p[i] = v; } +__device__ inline void Store(__hip_bfloat16* p, int64_t i, float v) { + p[i] = __float2bfloat16(v); // round-to-nearest-even, as the CUDA path and the host do +} + +// One block per token. Byte-for-byte the CUDA donor: same reduction tree, same +// scale loop, same cache offsets, same neox/gpt-j pairing, and the same EARLY +// RETURN on an out-of-range position — which happens AFTER the latent half has +// been written, so an out-of-range token still gets its norm and keeps its +// pe_out row untouched, exactly as RopeFromCache's own `continue` does. +template +__global__ void FusedNormRopeKernel(T* latent_out, T* pe_out, const T* x, const T* w, + const Tid* positions, const T* cache, int64_t cache_rows, + int64_t off, int rot, int64_t half, int64_t x_row_stride, + int64_t lat_row_stride, int64_t pe_row_stride, + bool is_neox_style, bool gemma, float eps) { + const int64_t row = blockIdx.x; + const T* xrow = x + row * x_row_stride; + T* lrow = latent_out + row * lat_row_stride; + T* prow = pe_out + row * pe_row_stride; + + // --- latent RMSNorm over [0, off) — identical to RmsNormRowKernel. ---------- + __shared__ float partial[kBlock]; + float acc = 0.0f; + for (int64_t j = threadIdx.x; j < off; j += kBlock) { + const float v = Load(xrow, j); + acc += v * v; + } + partial[threadIdx.x] = acc; + __syncthreads(); + for (int s = kBlock / 2; s > 0; s /= 2) { + if (static_cast(threadIdx.x) < s) partial[threadIdx.x] += partial[threadIdx.x + s]; + __syncthreads(); + } + const float inv = 1.0f / sqrtf(partial[0] / static_cast(off) + eps); + for (int64_t j = threadIdx.x; j < off; j += kBlock) { + float wj = Load(w, j); + if (gemma) wj += 1.0f; + Store(lrow, j, Load(xrow, j) * inv * wj); + } + + // --- decoupled-pe RopeFromCache over [off, off+rot) — identical to + // RopeFromCacheK (single vector, base rope, positions rank-1). ----------- + const int64_t position = static_cast(positions[row]); + if (position < 0 || position >= cache_rows) return; + const int64_t cache_offset = position * rot; + for (int64_t pair = threadIdx.x; pair < half; pair += kBlock) { + const float c = Load(cache, cache_offset + pair); + const float sn = Load(cache, cache_offset + half + pair); + const int64_t first = is_neox_style ? pair : pair * 2; + const int64_t second = is_neox_style ? pair + half : pair * 2 + 1; + const float xr = Load(xrow, off + first); + const float yr = Load(xrow, off + second); + Store(prow, first, xr * c - yr * sn); + Store(prow, second, xr * sn + yr * c); + } +} + +template +void LaunchFusedNormRope(hipStream_t stream, Tensor& latent_out, Tensor& pe_out, + const Tensor& x, const Tensor& w, const Tensor& positions, + const Tensor& cache, const RmsNormArgs& norm_args, + const RopeArgs& rope_args) { + const int64_t t = x.shape[0]; + const int64_t off = w.shape[0]; + const int rot = rope_args.rotary_dim; + const int64_t half = rot / 2; + if (t == 0) return; + const unsigned rows = static_cast(t); + if (positions.dtype == DType::kI32) { + FusedNormRopeKernel<<>>( + latent_out.Ptr(), pe_out.Ptr(), x.Ptr(), w.Ptr(), + positions.Ptr(), cache.Ptr(), cache.shape[0], off, rot, half, + x.stride[0], latent_out.stride[0], pe_out.stride[0], rope_args.is_neox_style, + norm_args.gemma, norm_args.eps); + } else { + FusedNormRopeKernel<<>>( + latent_out.Ptr(), pe_out.Ptr(), x.Ptr(), w.Ptr(), + positions.Ptr(), cache.Ptr(), cache.shape[0], off, rot, half, + x.stride[0], latent_out.stride[0], pe_out.stride[0], rope_args.is_neox_style, + norm_args.gemma, norm_args.eps); + } + Check(hipGetLastError(), "launch"); +} + +} // namespace + +// The registered FusedNormRopeFn (include/vt/ops.h:3388). Signature is the +// shared vt one, unchanged — that is the whole contract a backend has to meet. +void FusedNormRopeKernelRocm(Queue& q, Tensor& latent_out, Tensor& pe_out, const Tensor& x, + const Tensor& w, const Tensor& positions, const Tensor& cache, + const RmsNormArgs& norm_args, const RopeArgs& rope_args) { + EnsureQueueDevice(q); + switch (x.dtype) { + case DType::kF32: + LaunchFusedNormRope(AsStream(q), latent_out, pe_out, x, w, positions, cache, + norm_args, rope_args); + break; + case DType::kBF16: + LaunchFusedNormRope<__hip_bfloat16>(AsStream(q), latent_out, pe_out, x, w, positions, + cache, norm_args, rope_args); + break; + default: + VT_CHECK(false, "rocm fused_norm_rope: unsupported dtype (f32/bf16 only)"); + } +} + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_ops.hip b/src/vt/rocm/rocm_ops.hip index 29a1bd78a..9e6a72122 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -115,6 +115,13 @@ void SigmoidGateBf16KernelRocm(Queue& q, Tensor& out, const Tensor& attn, void Exl3GemmKernelRocm(Queue& q, Tensor& c, const Tensor& a, const Tensor& trellis, const Tensor& suh, const Tensor& svh, Tensor& a_had, const Exl3GemmArgs& args); +// ROCM-FUSED-NORM-ROPE (rocm_mla_fused_norm_rope.hip, #2564): the MLA +// norm-rope fold. Registering it is what makes `mla_attention.cpp`'s `fused_nr` +// predicate true on this backend, and that predicate -- not the op call -- is +// what a keep-quant MLA checkpoint depends on. +void FusedNormRopeKernelRocm(Queue& q, Tensor& latent_out, Tensor& pe_out, const Tensor& x, + const Tensor& w, const Tensor& positions, const Tensor& cache, + const RmsNormArgs& norm_args, const RopeArgs& rope_args); void AttnQkNormRopeGateKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& gate_out, const Tensor& qgate, const Tensor& kf, const Tensor& q_norm, const Tensor& k_norm, @@ -252,6 +259,9 @@ struct Registrar { static_cast(&SigmoidGateBf16KernelRocm))); RegisterOp(OpId::kExl3Gemm, DeviceType::kROCM, reinterpret_cast(static_cast(&Exl3GemmKernelRocm))); + RegisterOp( + OpId::kFusedNormRope, DeviceType::kROCM, + reinterpret_cast(static_cast(&FusedNormRopeKernelRocm))); RegisterOp(OpId::kAttnQkNormRopeGate, DeviceType::kROCM, reinterpret_cast( static_cast(&AttnQkNormRopeGateKernelRocm))); diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 4c596bfd5..8b86e9e42 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1919,6 +1919,108 @@ TEST_CASE("RmsNormGated and SigmoidGate match the CPU oracle") { } +TEST_CASE("FusedNormRope matches the CPU oracle within NMSE <= 5e-4, both styles") { + // The Tier-A2+A5 MLA norm-rope fold (ROCM-FUSED-NORM-ROPE, #2564). It is the + // ONE op in this harness whose ABSENCE on a backend is a refusal rather than + // a slowdown: `mla_attention.cpp` branches on + // `vt::OpRegistered(kFusedNormRope, device)` BEFORE calling it, and the + // fallback that branch selects row-slices a weight that a keep-quant MLA + // checkpoint stores block-quantized. So a device that skips this case cannot + // serve GLM-5.3 at all, and a SKIP here is a gap, not a pass. + // + // Geometry is DeepSeek/GLM-shaped but small: off = kv_lora_rank, rot = + // qk_rope_head_dim. `off` is deliberately NOT a multiple of the 256-wide + // block, so the strided reduction loop's tail is exercised rather than + // divided away. + constexpr int64_t kTokens = 13, kOff = 130, kRot = 16, kMaxPos = 64; + constexpr float kEps = 1e-6f; + + const std::vector x0 = RandomVec(kTokens * (kOff + kRot), 9101); + const std::vector w0 = RandomVec(kOff, 9102, -1.0f, 1.0f); + const std::vector cache = RandomVec(kMaxPos * kRot, 9103, -1.0f, 1.0f); + // Positions are NOT 0..n-1: a kernel reading the token index instead of the + // position passes on the identity mapping and fails here. + std::vector pos(kTokens); + for (int64_t i = 0; i < kTokens; ++i) { + pos[static_cast(i)] = int32_t((i * 5 + 2) % kMaxPos); + } + + // NeoX rotates (pair, pair+half); GPT-J style rotates (2*pair, 2*pair+1). + // DeepSeek-V2/V3 and GLM-5.3 use the GPT-J form, so `false` is the arm this + // model actually runs and `true` is the one a hardcoded kernel would break. + for (bool neox : {true, false}) { + CAPTURE(neox); + vt::RmsNormArgs na; + na.eps = kEps; + na.gemma = false; // DeepSeek/GLM: plain RMSNorm, not (1 + w) + vt::RopeArgs ra; + ra.rotary_dim = kRot; + ra.is_neox_style = neox; + + std::vector ref_lat(kTokens * kOff), ref_pe(kTokens * kRot); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector cx = x0, cw = w0, cc = cache; + std::vector cpos = pos; + Tensor tl = Tensor::Contiguous(ref_lat.data(), DType::kF32, cd, {kTokens, kOff}); + Tensor tp = Tensor::Contiguous(ref_pe.data(), DType::kF32, cd, {kTokens, kRot}); + Tensor tx = Tensor::Contiguous(cx.data(), DType::kF32, cd, {kTokens, kOff + kRot}); + Tensor tw = Tensor::Contiguous(cw.data(), DType::kF32, cd, {kOff}); + Tensor tpos = Tensor::Contiguous(cpos.data(), DType::kI32, cd, {kTokens}); + Tensor tc = Tensor::Contiguous(cc.data(), DType::kF32, cd, {kMaxPos, kRot}); + vt::FusedNormRope(cq, tl, tp, tx, tw, tpos, tc, na, ra); + cpu.DestroyQueue(cq); + } + + // The CPU oracle is itself the composite {RmsNorm ; RopeFromCache}, so an + // all-zero reference would make the NMSE ratio vacuous. Assert it is not. + double mag = 0.0; + for (float v : ref_lat) mag += std::fabs(static_cast(v)); + for (float v : ref_pe) mag += std::fabs(static_cast(v)); + REQUIRE(mag > 1.0); + + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kFusedNormRope, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + + DevBuf dx(dev, q, kTokens * (kOff + kRot)), dw(dev, q, kOff), + dc(dev, q, kMaxPos * kRot), dlat(dev, q, kTokens * kOff), + dpe(dev, q, kTokens * kRot); + dx.Upload(x0); + dw.Upload(w0); + dc.Upload(cache); + // pe_out is written only for an IN-RANGE position, so pre-seed both + // outputs with a value the kernel must overwrite. Zeros would let a + // kernel that wrote nothing pass wherever the oracle happened to be small. + dlat.Upload(std::vector(kTokens * kOff, -7.5f)); + dpe.Upload(std::vector(kTokens * kRot, -7.5f)); + void* dpos = dev.Alloc(kTokens * sizeof(int32_t)); + dev.Copy(q, dpos, pos.data(), kTokens * sizeof(int32_t)); + dev.Synchronize(q); + + Tensor tl = Tensor::Contiguous(dlat.ptr(), DType::kF32, d, {kTokens, kOff}); + Tensor tp = Tensor::Contiguous(dpe.ptr(), DType::kF32, d, {kTokens, kRot}); + Tensor tx = Tensor::Contiguous(dx.ptr(), DType::kF32, d, {kTokens, kOff + kRot}); + Tensor tw = Tensor::Contiguous(dw.ptr(), DType::kF32, d, {kOff}); + Tensor tpos = Tensor::Contiguous(dpos, DType::kI32, d, {kTokens}); + Tensor tc = Tensor::Contiguous(dc.ptr(), DType::kF32, d, {kMaxPos, kRot}); + vt::FusedNormRope(q, tl, tp, tx, tw, tpos, tc, na, ra); + dev.Synchronize(q); + + CHECK(Nmse(ref_lat, dlat.Download()) <= kNmseTol); + CHECK(Nmse(ref_pe, dpe.Download()) <= kNmseTol); + + dev.Free(dpos); + dev.DestroyQueue(q); + } + } +} + TEST_CASE("AttnQkNormRopeGate matches the CPU oracle within NMSE <= 5e-4") { // Fused full-attention preamble: split q|gate + (gemma) qk-RMSNorm(Dh) + // partial NeoX RoPE-from-cache + gate passthrough. Padded qgate/kf token