From 120456ac7a92fa589d24768ba8cff12576771538 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 2 Sep 2026 11:20:10 +0000 Subject: [PATCH 1/5] spec(BACKEND-ROCM): port kFusedNormRope to ROCm, because the branch that needs it asks a native-only question GLM-5.3 loads on `strix:gpu0` and its first forward throws (#2564). The MLA block's `fused_nr` predicate ends in `vt::OpRegistered(kFusedNormRope, device)`, which is a native-only probe by design, and ROCm registers no such op. With every environment variable unset the predicate is false, the split A-projection path row-slices `kv_a_proj_with_mqa`, and this checkpoint's q8_0 weight has no row slice. This spec takes the first of the three repairs #2564 prices, and says why it rejects the other two. Both halves of the composite are already native ROCm kernels -- `rocm_rmsnorm.hip`'s wavefront-agnostic tree reduce and `rocm_dense_basic.hip`'s cache-read rotation -- so the port is the composition CUDA already makes rather than new arithmetic. Widening `OpRegistered` to see the reference tier would instead change what "available" means for every backend and every op, and it keys naturally on a host-addressability predicate that answers backwards on this board: `gfx1151` reports `pageableMemoryAccess=0`, so the two host-addressability predicates disagree (#2515). The upstream anchors are read on the pin rather than copied from the briefing, and one of them was wrong: `GlmMoeDsaForCausalLM` lives at `deepseek_v2.py:1930` at `5559679229`, not in `deepseek_v32.py`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/specs/rocm-fused-norm-rope.md | 179 ++++++++++++++++++++++++++ 1 file changed, 179 insertions(+) create mode 100644 .agents/specs/rocm-fused-norm-rope.md diff --git a/.agents/specs/rocm-fused-norm-rope.md b/.agents/specs/rocm-fused-norm-rope.md new file mode 100644 index 000000000..2890576fb --- /dev/null +++ b/.agents/specs/rocm-fused-norm-rope.md @@ -0,0 +1,179 @@ +# 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`. 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 + +Recorded in `## Outcome` when the row reaches `DONE`. + +## 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. From 793f864d976e17fc97473fd8bb073f0580ddf74a Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 2 Sep 2026 11:23:16 +0000 Subject: [PATCH 2/5] feat(BACKEND-ROCM): register kFusedNormRope on ROCm, so a keep-quant MLA checkpoint stops taking the path it cannot survive GLM-5.3 loads on `strix:gpu0` and emits nothing: the first forward throws in the MLA block (#2564). The cause is not the op call, it is the BRANCH above it. `mla_attention.cpp`'s `fused_nr` ends in `vt::OpRegistered(kFusedNormRope, device)`, a native-only probe by design, and ROCm registered no such kernel. So with every environment variable unset the split A-projection path was taken, and it row-slices `kv_a_proj_with_mqa` -- which a q8_0 weight has no way to provide, because a quant block spans whole rows. That makes this op unlike the seven other MLA ops ROCm still lacks. Those fall through to the portable reference tier and merely run on the host. This one is consulted BEFORE the call, so its absence is a refusal. The kernel is the composition of two kernels this backend already had natively: `rocm_rmsnorm.hip`'s shared-memory tree reduce, which uses no warp-level primitive and is therefore wavefront-width agnostic, and `rocm_dense_basic.hip`'s cache-read rotation. `kBlock` stays 256 for the reason `rocm_rmsnorm.hip:41` already gives -- four whole wavefronts, and the same reduction ORDER as the CUDA and CPU siblings, which is what makes an NMSE bar mean anything. The refusal message no longer names one cause when there are two. It reads the predicate terms back: either the environment override is set, or this backend registers no native kernel and `OpRegistered` cannot see the reference tier. The previous text enumerated the backends that HAVE the op and forgot the ones that do not, so it sent a reader looking for a variable nobody had set. Gated by a new FusedNormRope arm in the cross-device harness, at NMSE <= 5e-4 against the CPU oracle in both rope styles, with `off` deliberately not a multiple of the block width and the outputs pre-seeded to a value the kernel must overwrite. Refs #2564 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- CMakeLists.txt | 2 + .../layers/attention/mla_attention.cpp | 33 +++- src/vt/rocm/rocm_mla_fused_norm_rope.hip | 181 ++++++++++++++++++ src/vt/rocm/rocm_ops.hip | 10 + tests/vt/test_backend_cross_device.cpp | 102 ++++++++++ 5 files changed, 321 insertions(+), 7 deletions(-) create mode 100644 src/vt/rocm/rocm_mla_fused_norm_rope.hip 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/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..95c2c507c 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 double kEps = 1e-6; + + 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 From d2e18fed559890ea920b6bc668cc23c3943cd05d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 2 Sep 2026 11:28:57 +0000 Subject: [PATCH 3/5] docs(BACKEND-ROCM): recount the op table with a scan that survives a wrapped argument list Found while adding the `kFusedNormRope` registration this branch carries. The new call wraps after `RegisterOp(`, the way eight of its neighbours already do, and the recount command this document prescribes did not see it -- the count read the same before and after adding an operation. `grep` is line-based, so `RegisterOp(OpId::[A-Za-z0-9_]*` never matches a call whose argument list breaks after the paren. A contributor who runs the documented command and quotes its answer quotes an undercount, and the command cannot detect that it undercounted. The recorded 44 beside it was stale by eight regardless. The replacement reads the file whole with `-z` and names the `DeviceType` in the pattern, because the old one also counted every backend's registrations in the directory rather than the one a reader asked about. The cell now says what the table IS. `AGENTS.md` forbids storing a measurement of one file inside another, and a live count in a projection document is exactly that shape -- so the fix is to remove the number, not to correct it. Closes #2573 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- docs/ROCM.md | 16 +++++++++++++--- 1 file changed, 13 insertions(+), 3 deletions(-) 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 | From b413e323be50822dcbecb30bd61dc90333a416b5 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 2 Sep 2026 11:30:41 +0000 Subject: [PATCH 4/5] test(BACKEND-ROCM): give the FusedNormRope case's eps the type the field has `RmsNormArgs::eps` is a `float`. The case declared its constant `double` and assigned it, which is a narrowing conversion the compiler is free to warn about under a stricter flag set than this tree carries today, and which states the wrong thing about the value either way: the number that reaches the kernel is the float, so that is the number the test should name. No behaviour changes -- 1e-6 is exactly representable either way. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- tests/vt/test_backend_cross_device.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 95c2c507c..8b86e9e42 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1933,7 +1933,7 @@ TEST_CASE("FusedNormRope matches the CPU oracle within NMSE <= 5e-4, both styles // 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 double kEps = 1e-6; + 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); From 67d2966f4b2b6f0a5bf54c659eb3183c68d98ca7 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Wed, 2 Sep 2026 13:59:41 +0000 Subject: [PATCH 5/5] record(BACKEND-ROCM): GLM-5.3 generates on gfx1151, and five ops that serve it are on the host The checkpoint row said this artifact loads and emits nothing. It now emits ` Paris, which is` on `strix:gpu0` through the shipped entry point, so the row says that instead, with the five reference-tier ops beside it -- because configuration alone cannot tell a device run from a host fallback, which is the failure #2505 names. No speed number is offered. `docs/ROCM.md` disqualifies a performance result from a run with a non-zero reference-tier hit count, and this run has five. The 3516.719 s the harness printed is recorded as the cost of a host-tier MLA arm and explicitly not as throughput. The spec's evidence section carries the measurements: the build, both assertion counts on the focused gate, the CPU-only control that makes those counts mean something, the op-id mapping cross-checked against the run's own named lines, and the mutation ladder with its restore proof. One instrument is recorded as useless rather than quietly dropped. Grepping a passing doctest run for the device name returns zero, because `CAPTURE` prints only on failure. The assertion count is what discriminates a device arm that ran from one that skipped: 20 on the board, 2 on a CPU-only build of the same tree. Refs #2564 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/specs/rocm-fused-norm-rope.md | 70 ++++++++++++++++++++++++++- docs/USAGE.md | 2 +- 2 files changed, 69 insertions(+), 3 deletions(-) diff --git a/.agents/specs/rocm-fused-norm-rope.md b/.agents/specs/rocm-fused-norm-rope.md index 2890576fb..8bc9a92f7 100644 --- a/.agents/specs/rocm-fused-norm-rope.md +++ b/.agents/specs/rocm-fused-norm-rope.md @@ -5,7 +5,7 @@ Issue: [#2564](https://github.com/mudler/vllm.cpp/issues/2564) ## Now -`ACTIVE`. Base `4fe3852b119e40cda05c0fbcf64d3e2a4796ada2`. +`ACTIVE`, measured. Base `4fe3852b119e40cda05c0fbcf64d3e2a4796ada2`. ## Scope @@ -156,7 +156,73 @@ inspection. ## Evidence -Recorded in `## Outcome` when the row reaches `DONE`. +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 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 |