feat(rocm): f32-query decode GQA kernel with bf16 and fp8 KV cache support - #8
Closed
ghazni101 wants to merge 13 commits into
Closed
feat(rocm): f32-query decode GQA kernel with bf16 and fp8 KV cache support#8ghazni101 wants to merge 13 commits into
ghazni101 wants to merge 13 commits into
Conversation
…Q5_K GEMV before tuning it (mudler#2250) (mudler#2255) Opens W11's lever #2 as its own row, and narrows it, because reading the prior work first changes what the lever is. W11 measured the keep-quant GEMVs at ~87% of Laguna decode GPU and called the lever "BW-tuning". This spec does not inherit that label. The sibling kernel carrying the other 24.7% was measured at the counter and is memory-LATENCY- and LSU-pipe-bound rather than bandwidth-bound: long_scoreboard 54-57 at 72-75% occupancy, L1 hit 96.6% so the 16x sector over-fetch never reaches DRAM, lg_throttle 74.2 on the weight unpack, and local_ld/local_st both zero, which refuted the register-spill hypothesis outright. Five structural levers there came back flat or refuted, including multi-row/prefetch ILP re-confirmed as a wash, and a MEASURED FLOOR is recorded. That half is CLOSED and this row says so, because re-opening it on a bandwidth premise would repeat five bricks of known-dead work. What IS open is the grouped Q4_K/Q5_K kernel, which is a different kernel and the bigger share at 62.1%. Its own spec names the deferral rather than this one inferring it: "MMVQ warp-per-output is the correctness-first structure; tensor-core tiling is a later speed brick." One warp per output, lanes striding the K super-blocks, never bandwidth- or ILP-tuned on CUDA. SO W1 IS A MEASUREMENT AND NOT AN OPTIMISATION, and the spec is written to make choosing a lever from the label impossible. It enumerates four readings the counters could give — bandwidth, latency, LSU-pipe, occupancy — and names a different next wave for each, including that a latency reading means the Q8_0 kernel's five refuted levers are the prior and the expected value of retrying them is low. A null result is recorded as the finding: two kernels sharing a floor is more useful than a sixth refuted brick. It also fixes the trap that has already produced one wrong attribution in this tree. The profile must target DECODE by kernel name past prefill, because a whole-run aggregate folds in prefill and one-time load-path work — which is how a contaminated kern_sum once yielded a "kernels already at parity" claim that a clean graph A/B later reversed. No llama.cpp denominator is quoted, and W11's own "~22% of peak vs llama.cpp ~76%" is explicitly marked as inheriting the mudler#1003 supersession, so it cannot be used as a target either. Gates: `check-agent-record` and `check-conflict-markers` rc=0. No code. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [Claude Code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…e blocker was a frame error and the row is SPIKE on a committed plan (mudler#2251) GLM-5.3 is 97.49% routed experts, so the hardware blocker was a frame error and the row moves to SPIKE on a committed plan mudler#2194 concluded that `GlmMoeDsaForCausalLM` cannot run on any fleet device, and it sized RESIDENT weights to get there. That is the right frame for a dense model and the wrong one for a 256-expert MoE. This change writes the port plan as section 3 of `.agents/specs/glm-dsa-latest-deepseek.md`, moves `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` from `BLOCKED` to `SPIKE`, and carries the records that move with it. No product code, no pin advance, no build, no GPU lease, no download. ## The arithmetic, recomputed rather than copied Recomputed from `zai-org/GLM-5.3`'s own `config.json` and checked against the checkpoint's `model.safetensors.index.json` (`metadata.total_size = 755,617,140,416` over 118,629 tensors): the routed experts are 734,439,407,616 of 753,328,717,824 parameters -- 97.49% streamable, 2.51% resident -- and the total reproduces the API's measured `safetensors.total` to -1,222,656 params, or -0.00016%. That is tighter than mudler#2214's own 745.8B model, which omitted the MTP block's 256 experts and therefore put resident 2.1B too high; bf16 resident is 35.18 GiB, not 39.19. Both figures favour the argument, so the correction does not change the verdict, and it is stated because a record written from another document's figures inherits its errors. ## The number that actually decides the row is measured from the artifact `unsloth/GLM-5.3-GGUF` held one arm at 319.41 GiB on 2026-08-28. At revision `346b3591c7f28d1a23716f97a065ecf12ec14771` it holds twelve arms and 5542.40 GiB. A full GGUF header census of `UD-IQ1_S` -- 6 shards, 1809 tensors matching the file's own `split.tensors.count`, read by HTTP range request at a cost of about 9.6 MB with nothing downloaded -- splits it into 228 `*_exps.weight` towers at 187.312 GiB and 1581 resident tensors at 14.511 GiB. The resident class is about 14.5 GiB in every published UD arm, because the recipe keeps every non-expert tensor at Q4_K or better, so the residency plan is arm-independent. One decode step at c=1 touches 75*3*8 = 1800 distinct slices, 11.21 GiB of uniform slots, so resident plus a 4096-slot cache is 40.01 GiB against 119.631 GiB on `dgx:gpu0`. A `UD-*` name is a target average and not a format, and the census is what proves it: `UD-IQ1_S` is 106 IQ1_S + 71 IQ3_XXS + 44 IQ2_XXS + 4 IQ4_XS + 3 K-quant expert tensors, and `UD-Q2_K_XL` holds exactly two Q2_K tensors out of 1809. ## The verdict changed under this branch, and the second commit is why The census was measured at the base `60a6dd97b`, where `IQ4_XS` (ggml id 23) and `IQ2_XS` (17) had no vt block dtype and no decoder, so the arm refused at load with a message naming the type. Merging `origin/main` brought in `94de63ff5` (mudler#2245), landed for the sibling GLM-5.3-Flash row two hours earlier, which ports BOTH dequantizers and NEITHER keep-quant `vec_dot`. `gguf_keep_quant.cpp::KeepQuantDType` returns false unless `vt::cpu::HasQuantDotKernel(dt)`, so both new types now EXPAND TO bf16 at load. And `gguf_device_fit.cpp:85-100` walks every `*_exps.weight` tensor and returns false for the WHOLE arm the moment one is not `kKeepQuant` or `kKeepF16`, so four offending tensors out of 228 drop the entire model out of the expert-streaming lane. One `*_exps` tower is 3,221,225,472 elements, 6.000 GiB at bf16: `UD-IQ1_S`'s four IQ4_XS towers go 6.375 -> 24.000 GiB, `UD-Q2_K_XL`'s 148 IQ2_XS towers go 128.344 -> 888.000 GiB, and the uniform slot grows from 6.375 MiB to a 24.00 MiB bf16 slice, making a 4096-slot cache 96.00 GiB. So the row is blocked on exactly one kernel and it is a `vec_dot` rather than a decoder: `VecDotIQ4_XSQ8_K`, four tensors `blk.{8,75,76,77}.ffn_down_exps.weight`, which also unlocks `UD-IQ2_M`. `QUANT-GGUF-IQ4_XS` already owns it. The general defect is named and not repaired here: landing a decoder without its `vec_dot` converts a loud refusal into a silent 3.4x memory multiplier that no token gate can see. ## One premise of mudler#2194 is corrected The pinned vLLM class CAN load this checkpoint. It never reads `indexer_types` -- zero occurrences tree-wide at `555967922` -- it DERIVES the schedule at `deepseek_v2.py:1092-1103` from `index_topk_freq` and `index_skip_topk_offset`, and drops surplus checkpoint indexer weights at `:1566-1582`. The config's `indexer_types`, that derivation, and llama.cpp `b10451`'s hardcoded `GLM_5_2_DEFAULT_INDEXER_TYPES` (`src/models/glm-dsa.cpp:6-27`) agree bit for bit over all 78 layers, and the checkpoint ships `self_attn.indexer.*` on exactly 22 of 79 blocks. `modules_to_not_convert`'s `self_attn.indexers_proj` matches no shipped tensor and no upstream module at the pin, so it is a config-level shorthand and a loader must not mirror it. ## What the streamer actually provides, and the delta `expert_streamer.cpp` is real and not turnkey. Its wiring lives entirely inside `qwen3_5.cpp`, which is the only model translation unit that constructs it; `deepseek_v2.cpp` has zero references. The default slot budget is 64 against a 1800-slice working set, and below the working set the model does not fail, it silently reads the mmap. There is no prefetch and no async I/O, eviction is an O(resident) linear scan per miss, and no device destination is wired -- which is why this port is GB10-shaped: `qwen3_5.cpp:6199` takes the slot arm only for `cpu || host_memory_is_device_addressable()`. Against `DeepseekV2ForCausalLM` the delta is smaller than this spike's own DSA verdict implies, because a device-native DSA lightning indexer now lives in the shared MLA block (`mla_attention.cpp:598-745`, CPU+CUDA `DsaIndexerLogits` / `DsaTopkSelect`, reached in production by `Dots3NoteForCausalLM`) and this model's MLA geometry already validates and dispatches to native kernels. Net-new: the indexer KV side cache (mudler#1925), sparse prefill, the heterogeneous per-layer schedule with `skip_topk` selection reuse, the fp32 router GEMM, a `"glm-dsa"` `kGgufArchArms` row, and lifting the streaming seam out of `qwen3_5.cpp`. ## The gate is the honest cost No end-to-end token gate against vLLM is reachable on this fleet. vLLM implements the architecture and cannot run 703.74 GiB on a 119.631 GiB unified device whose host RAM is the same pool, and it has no GGUF path for this architecture. Four gates are reachable and named: module parity against the pin on CPU, a headers-only structural loader gate, a streamed-vs-resident identical-logits gate that needs no oracle, and an llama.cpp `b10451` floor on the identical artifact. Speed is an open gap by construction. Eight waves are planned with scope, exclusions, anchors, tests and gates; none has landed. Twelve `## Owed` entries record what is not settled and what would discharge each. ## Records `RUNNABLE_BASELINE` in `scripts/check-gate-commands.py` is re-pinned in this change because the row leaves the gate-obliged population, not because it lost a command, with a dated note naming the row and the reason. The `### Gates` section that earned the credit is untouched, and section 3.6 adds four more. The matrix rollup moves BLOCKED 5 -> 4 and SPIKE 9 -> 10, the checklist mark moves from `🚫` to `📋`, `.agents/claims/CLAIM-MODEL-GLM-MOE-DSA.md` is new, and one row is appended to `.agents/issue-index.md`. `scripts/agent-preflight.sh --fail-on-skip`: All gates green, zero skips. Closes mudler#2214 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…ur render is no worse than the oracle on blockiness (mudler#1854, mudler#2220) (mudler#2210) measure(LTX25-ORACLE-ABSOLUTE): mudler#1854's reading is taken, and our render is no worse than the oracle on blockiness (mudler#1854, mudler#2220) mudler#1854 was filed rather than closed with a proxy, on the reasoning that "a proxy for perceptual quality that measures nothing is worse than a declared gap". The gate it asked for landed in mudler#2134 with no reading behind it. This change takes the reading. `rc` job `4b0666ee-248c-45fc-9de6-372b6d0c1fab` on `dgx:gpu0` rendered the manifest's exact request -- 320x192, 25 frames, 8 steps, seed 42, 240 video tokens -- on the bf16 arm the oracle used, and the absolute comparison returned PASS against both reference forms. blockiness_grid8 ours 1.022135 <= 1.143393 margin +0.121257 blockiness_grid32 ours 1.025445 <= 1.148672 margin +0.123227 both _defined 0 of 1600 bands collapsed READING NO_WORSE_THAN_ORACLE_ON_BLOCKINESS VERDICT PASS (exit 0) The bound was recomputed, not transcribed: the JSON carries `reference/bounds/blockiness_grid8/frame_max = 1.1433929206406797` and `digests_verified = 25`. The 25 NAS PPM frames and the committed mp4 both returned exit 0, which re-runs the spec's claim that the two forms agree on the gated bound instead of leaving it a number somebody wrote down. ## Three blockers cleared to get there, each located rather than guessed **The bf16 text tower.** mudler#2140's refusal was re-run on the bytes that produced it and is gone: `in_features` resolves to 188160 rather than the doubled 376320, `quantized_modules = 0`. **The 42 GB bf16 DiT, which nothing had ever asked about.** Every LTX-2.5 render in this tree had loaded the NVFP4 or FP8 transformer. `scripts/probe_ltx2_dit_load.cpp` is new and answers it in 1.8 s off the header: arm `kNone`, all 4091 contract tensors present at the contract's shape and byte count, 258 unbound tensors that are exactly the two `*_embeddings_connector` families `UnportedFamilies` skips, so no `allow_unported_modules`. It is a probe and not a load, deliberately, and it prints that it establishes no render before its `OK`; it is also shown able to say no, exiting 1 on a real refusal when pointed at the text encoder. **[mudler#2220](mudler#2220), a defect in this row's own harness, found by execution.** A CUDA toolkit staged off CIFS carries no symlinks, and the reconstruction used `${f#*.so.}`, which strips the SHORTEST prefix and yields `13.3.29` rather than `13`. The second `ln` therefore linked `libcudart.so.13.3.29` to ITSELF and never created `libcudart.so.13`, the SONAME `ld` resolves versioned undefined symbols against. CMake reported `Found CUDAToolkit`, every CUDA TU compiled, and the job died 21 minutes later with 38 `undefined reference to ...@libcudart.so.13`. `need_ok` tested `libcublasLt.so`, the one link the loop did create correctly, so the precondition passed on a toolkit that could not be linked against. It was latent, not new: the staging branch is a FALLBACK and every earlier lease found `/usr/local/cuda`. The two runs A/B in their own configure logs -- `/usr/local/cuda` 13.0.88 built in 1192 s, staged `/root/cudatk` 13.3.73 failed at link -- and `dgx:gpu0` went `unhealthy ... worker_lost` for 3h20m between them and returned without a toolkit. Fixed by taking the MAJOR, preferring `ldconfig -n` so the name comes from each object's own `DT_SONAME` rather than string surgery on a filename, and asserting the postcondition in seconds instead of after a 21-minute build. Red-before/green-after on a replica of the CIFS layout: the old loop creates no `.so.13`, the new logic creates both, and the guard FAILS on the old layout, PASSES on the new, and FAILS on the real NAS source -- so it discriminates rather than passing by construction. ## `--steps` is proven by execution This was the row's one wired-but-never-run path. `steps_requested=8 steps_observed={8} dit_forwards=32`, where the observed set is the distinct denominators of the sampler's own `step k/M` lines and M is `sigmas.size() - 1`. A number the sampler computed, not the flag echoed back; 32 forwards over 8 steps is 4 per step, the guided denoiser's quartet, so the count corroborates the schedule rather than merely agreeing with it. The silent failure mode was ruled out beforehand: a phase carrying its own sigmas with `allow_request_sigmas` true would have kept its schedule, ignored the override, put a 30-step render against an 8-step reference, and passed. ## What the green does NOT say We are less blocky than the reference's own MEAN, not merely under its maximum. But on the two REPORTED statistics we sit outside its per-frame range in the same direction: sharpness 10.5176 against a reference minimum of 10.8391, clipped fraction 0.000758 against 0.001226. Less blocky, less sharp and less clipped is one coherent picture -- our render is somewhat SMOOTHER than upstream's -- and smoothness is exactly what a one-sided blockiness ceiling cannot see. Neither is gated, and section 5 gives the measured reason no bound can be derived for either. It is recorded beside the green because a reader who takes PASS as "matches upstream" would be wrong. The gate's claim is its name. Prompt adherence is still not measured, here or anywhere in this tree. **mudler#1854 is NOT closed by this change.** ## Evidence C0 on our render: 25 distinct frame hashes of 25, zero near-uniform frames, minimum per-frame variance 2186.296, zero zero-motion pairs, mean adjacent MAD 5.4060. All four checkpoint sha256 recomputed inside the lease on the staged copies, all four matching the manifest -- a second independent reading after the CIFS-side verification. The CUDA unit gate ran BEFORE the render at 23 cases / 806 assertions / 0 failed, checked rather than assumed, because a doctest binary that skips everything also exits 0. Binary `7b1f4367...`, library `9e3dc6f4...`, harness `5649b4e8...`, source `0002ddfba`. `docs/USAGE.md` carried two claims this run falsified and both are corrected: the bf16 tower row said "no render has yet been gated on this one", and the DiT row named no gate at all. `.agents/environment.md` gains the SONAME fact, because a staged CUDA runtime whose links did not survive CIFS will bite the next row. Closes mudler#2220, whose fix is verified by this very render: the build that produced these artefacts is the first this harness has ever completed on a staged toolkit, and the lease recorded the rebuilt SONAMEs (`libcudart.so.13 -> libcudart.so.13.3.29`) in its own PROVENANCE. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…e, and the review repairs that the first attempt at saying so still owed (mudler#2218) (mudler#2254) fix(MODEL-MM-QWEN4-EXP): one gamma polarity for the whole architecture, and the review repairs that the first attempt at saying so still owed (mudler#2218) `Qwen4ExpWeights` holds every norm gamma in the RAW HuggingFace parameterization: `LoadNormBf16(..., unshift=true)` inverts the `+1` that ggml-org/llama.cpp#27742 bakes in at convert time, with `linear_attn.norm.weight` the one tensor that converter never folds. `vt::Qwen4ExpGatedResidual` documented the opposite — "hc_norm_w is vLLM's parameterization, i.e. ALREADY `1 + w_hf` ... This op never adds 1" — so the layer loop W5b is about to write would have handed it a gamma centred on zero and multiplied every hyper-connection norm by ~0. A plausible tensor, never a crash, and unreachable by any gate in this tree: the loader was right about its output and the op was right about its input, and no suite composed them. THE OP NOW ADDS THE 1. That direction rather than the other one, because it is the direction three of the four consumers had already taken and the count is checkable. `RunQwen4ExpQsaBlock` normalizes THREE gammas through `vt::RmsNorm(..., gemma=true)`, which is `out * (1 + w)` — `idx_q_norm` (`qwen4_exp_qsa_block.cpp:383-384`), `q_norm` (`:425-426`) and `k_norm` (`:441-442`). The fourth, `idx_k_norm`, never reaches `vt::RmsNorm`: it goes to `Qwen4ExpQsaIndex` (`:401-403`) and is consumed inside it by `vt::Qwen4ExpQsaCompress` (`:181`), which documents its `k_norm_weight` as "the HuggingFace gamma, applied as `(1.0 + weight)`, ... NOT vLLM's `out * weight`". The PLE host reference spells `(1.0 + weight)` inline at `qwen4_exp_ple.cpp:72`. Same polarity at all four, one of them through a different op — the earlier text said `vt::RmsNorm` for all four and cited three line pairs, and the fresh review counted the citations against the claim. It is also upstream verbatim: `Qwen4ExpTextRMSNorm.forward` is `output * (1.0 + self.weight.float())` over a zero-initialised parameter (transformers v5.16.0 `modeling_qwen4_exp.py:173-178`). So one rule now covers the model — every gamma raw, every consumer adds the 1 — instead of three tensors on one convention and a fourth on another. This is a deliberate reversal of the remedy mudler#2218 proposed. That issue asked the layer loop to fold `hc_norm`, `norm_key`, `norm_query` and `norm_conv` through `HcNormWeightFromHf` before use; folding the last three would have been the same defect moved one tensor to the left, because their consumers already add the 1. Only `hc_norm` ever disagreed. A further correcting comment is on the issue. ## The gate `tests/vllm/models/test_qwen4_exp_forward.cpp`, and what makes it an instrument rather than a restatement is that it LOADS the gamma. It drives the production `ModelRegistry::Load` over the synthetic `qwen4exp` file, takes `layers[0].attn_hc`, `layers[0].mlp_hc` and the model-level `use_combine` mixer straight out of `Qwen4ExpWeights`, runs them through the device op, and compares against the W3 host reference driven with the value the FILE carried. A test that builds its own gamma has already chosen the answer it then asserts. **It is the first `qwen4_exp` suite to compose the loader with an op, and NOT the first in this tree.** The earlier wording claimed the latter and was wrong: `tests/vllm/models/test_nemotron_h_paged_forward.cpp` and `tests/vllm/models/test_kimi_linear_paged.cpp` each call `ModelRegistry::Load` inside a `TEST_CASE` and drive the loaded weights, gammas included, through the device ops of a forward. The scoped claim is the one the argument needed: no `qwen4_exp` suite had put the two halves in one process, which is why eleven single-sided waves OF THIS ROW could not see the contradiction. Two mutations, each applied under a sha256 before/after proof with the build rc read before any test result and the tree restored byte-for-byte: M-P1 drop the `1 +` from the kernel build rc=0, forward 1/1 FAILED, hc_device 4/9 FAILED M-P2 loader stops unshifting `hc_norm` build rc=0, forward 1/1 FAILED at its precondition, gguf_weights 1/11 FAILED M-P2 is the half that matters most: it proves the case gates the LOADER side and not only the op, so a future edit that moves the fold back into `load_weights` cannot land silently. The injection arm of the new case is asserted SATURATED rather than counted as evidence. `2 * sigmoid(inject . normed / hc)` reaches 2.0 under both polarities at this fixture, so it discriminates nothing here; `mixed` is the discriminator and the case says so. The synthetic file moves to `tests/support/qwen4_exp_gguf_fixture.h` verbatim so the loader suite and this one share ONE builder. Its dimensions were each chosen to make a specific defect expressible — `kNumKHeads` 2 against `kNumVHeads` 6 so the V-head permutation is not its own inverse, `kPleRow` 96 so `ple_embed_dim` is distinct from both `hidden_size` and `hidden_size * ngram_heads`, gammas on a bf16-exact `1 + k/128` grid so the fold cannot round away. The loader suite is unchanged in behaviour: 11 cases and 2975 assertions before and after. ## The fold is f32, and it had drifted Before this change the wide-accumulator case in `test_qwen4_exp_hc_device.cpp` handed ONE identical `float` multiplier to both arms. The first version of this change left the kernel folding `1.0f + w` in f32 while the double reference folded `1.0 + (double)w`, so the case's own comment — "the only thing this widens is the reduction" — stopped being true. Nothing failed, which is the point. Measured on exactly that data by forcing the bound to `1e-30` and reading the logged worst deviation: reference folds `1.0f + w_hf`, widens after 1.17323e-06 reference folds `1.0 + (double)w_hf` 9.8457e-07 Both are far inside the band — the bound is `1e-5` and the `float ss` mutant reads 6.702e-4 — so no tolerance was ever at risk. The MEANING of the number was: 1.173e-06 is the value this file and the W5b-2 table record as "ours, double accumulator", and the drifted form no longer reproduced it. The reference now folds in `float` and widens afterwards, which restores the recorded figure exactly and leaves the reduction as the only widened thing. f32 is upstream's width, not a convenience: the Python `1.0` in `output * (1.0 + self.weight.float())` is a weak scalar, so the promotion stays fp32. AGENTS.md "Inherit vLLM defaults" decides it independently — f32 is the default and the wider value would have been the annotated exception, unannotated. ## The rest of the review's findings - **`qwen4_exp_registry.cpp` claimed the forward "is the only thing that does" refuse.** Two more do, both load-time: the safetensors arm refuses unconditionally, and the GGUF arm refuses a source that names the kind without carrying a file. The paragraph now counts three and points at their sites. That is the same class of overbroad claim this change exists to remove. - **A dead build define.** `tests/CMakeLists.txt` granted `test_qwen4_exp_forward` a `QWEN4_EXP_CKPT_FIXTURE_DIR` whose comment named "the reachability case". There is no such case: the file has one `TEST_CASE`, and neither it nor `tests/support/qwen4_exp_gguf_fixture.h` reads the define. Removed, with the comment now saying why no fixture dir is needed — the suite builds its `qwen4exp` file in memory. - **A latent name collision, recorded at both ends.** Two functions named `GroupedRmsNorm` now live in `namespace vllm::qwen4_exp` with OPPOSITE polarity: the public one (`qwen4_exp_hc.cpp:52`, `out * w`, fed through `HcNormWeightFromHf`) and a file-static one (`qwen4_exp_ple.cpp:55`, `out * (1 + w)`). Nothing collides — different signatures, and `qwen4_exp_ple.cpp` does not include `qwen4_exp_hc.h` — but a name shared across two gamma conventions is this issue's exact shape. Both sites now name the other and its polarity. Deliberately NOT renamed: the duplicate disappears when PLE's three norms move onto the standalone grouped-norm `vt::` op that `## Owed` item 1 already carries. - **Two merges of `origin/main`.** The branch was five commits behind, so both commit-walk gates SKIPPED and `--fail-on-skip` read rc 1; `origin/main` then moved again under the first rerun. Both merges carry an authored message with the trailer block, because a default merge commit has none and would red the gate the merge exists to un-skip, with no non-force repair once it lands. Neither merge shares a single file with this branch. ## What this does NOT do **`vt::Qwen4ExpGatedResidual` LANDS UNREACHED, by AGENTS.md "Nothing lands dead".** It and `vt::Qwen4ExpGatedResidualWriteBack` have no production call site at this merge commit; their only callers are `tests/vllm/models/test_qwen4_exp_hc_device.cpp` and the new `tests/vllm/models/test_qwen4_exp_forward.cpp`. That second suite reaches the production LOADER, which is what makes the fix gateable at all, but a test driving a loader is still a test, and reaching the loader does not reach the op. `Qwen4ExpTextModel::Forward` does not exist and `ForwardQwen4ExpForConditionalGeneration` still refuses by name before any downcast. **The wiring is owned by row `MODEL-MM-QWEN4-EXP` and by W5b under mudler#2031**, tracked by campaign mudler#1978, and the spec lists it under `## Owed`. Nothing here decodes a token. The refusal in `qwen4_exp_registry.cpp` now names the five measured prerequisites a layer loop actually lacks — a standalone grouped RMS norm for PLE, a paged QSA consumer, the group-2 block table, a MoE weight adapter and an externally linked mRoPE builder — instead of naming W2, W3 and W4, which landed. `.agents/specs/qwen4-exp-flash-next.md` carries each with its citation, and corrects two of its own earlier claims in place rather than deleting them. Nine suites, rebuilt and rerun at the head, build rc read before every result: forward 1/421, hc_device 9/87, hc 15/246, gguf_weights 11/2975, ple_device 10/538, qsa_device 12/4697, qsa_block 8/2831, kv_cache 4/399, scaffold 12/296. The eight pre-existing counts are identical across the op's contract change, which is the check that it is a re-parameterization and not a new answer. Closes mudler#2218. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…into a replayed graph (mudler#2252) (mudler#2271) fix(SPEC-DFLASH2): the capture-safe bound was a per-STEP value baked into a replayed graph (mudler#2252) `5f8a70705` stopped the draft's paged attention from calling `cudaStreamSynchronize` inside a CUDA graph capture. It set pa.max_seq_len = ctx_len + tq which is exact at capture and STALE on every replay after it, because the graph is captured once and replayed as the context grows. `MakeDeviceKVStore` states the invariant that makes capture legal at all, and it is the one this broke: the persistent buffers never move, so "a captured graph reads the growing context purely through the in-place `seq_lens` value". A host value derived from the current `ctx_len` is exactly what that forbids. So the first fix traded a LOUD failure for a SILENT one — an out-of-bounds read instead of a refused capture. That is the worse trade, and it is worth naming. MEASURED on `dgx:gpu0`, one boot, three arms, at the smallest workload that reproduces it (`max_num_seqs=1`, c=1, 64 tokens, k=7): A production default exit 134, IMA at `cudaMemcpyAsync` C A + CUDA_LAUNCH_BLOCKING=1 exit 134, IMA at `cudaGraphLaunch` B VT_DFLASH_PAGED=0 exit 0 C names the replay rather than the `cudaFree` that merely noticed the fault later, and B — which bypasses the paged route, the `P == 1` capture lane and every line the first fix touched — exits clean on the same binary. Together they place the fault inside the replayed graph and nowhere else. The bound now comes from the POOL: `pool_k.shape[0] * pool_k.shape[1]`, pages times page rows. It is replay-stable, and the read can never pass it because the runner refuses a request whose `ctx_len + append + (1+k)` would exceed the store's capacity. An upper bound is explicitly safe there — it only sizes grids and rounded dims, while per-request geometry stays on the DEVICE values (`ops.h:1551-1553`). THE GATE IS ON THE PROPERTY, NOT THE VALUE. A single-point assertion passes `ctx_len + tq` happily, which is why the first round's test did not catch this. The new case asserts the bound COVERS the worst-case replay sequence for every capture-time context in {0, 16, 1200} — the comparison the GPU reported as an illegal access. It also drops a tautology the first draft of it contained (`a.max_seq_len == b.max_seq_len` over identical arguments). Red-first, and the first attempt at that proof is recorded because it lied: mutating `max_seq_len` to a capture-time bound FAILED THE BUILD on `-Werror=unused-parameter`, so the stale binary ran and printed 16/16 PASSING. Redone with `(void)pool_capacity`, it reads `BUILD rc=0 / TEST rc=1` failing on `hm.max_seq_len >= worst_case_replay_seq`. All 29 `dflash` binaries are green. Found by the committed gate, which refused (`GATE_RC=2`, `RESULT_PRESENT=no`) rather than reporting a number. The same run produced a clean oracle denominator — 16.347 tok/s at 0.795% SM-clock spread, persistence enabled, no throttle reasons — so what is still owed is our own side of that comparison. No second index row: mudler#2252 already has one, and `check-agent-record` correctly refused the duplicate ("under `merge=union` a duplicate is what two branches appending the same issue look like"). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…udler#2244) The W5 wave needed its issue and its committed spec section before any implementation (AGENTS issue-first + spec-before-code). W3's leftovers are also off the owed list: they landed 2026-08-29 via mudler#2217, and ## Now still told a reader they were next. Three record moves, one unit: - .agents/issue-index.md: the mudler#2244 row — the staging upload still pays tt-metal's full creation path every step (UploadRowsBf16 builds a new ttnn::Tensor via from_span per upload), and the W4 record's named lever is the per-slot persistent device buffer written through the mesh command queue, with the tt-metal-internal half as a proof obligation. - specs/tenstorrent-qwen35.md ## Now: W3 leftovers recorded as landed (mudler#2201 via mudler#2217, a456e6e, suite at 44 cases / 4340 assertions); the owed-next pointer moves to W5. - specs/tenstorrent-qwen35.md Work breakdown: the W5 section — per-slot persistent buffer, MeshCommandQueue::enqueue_write(_shard) writes, the mudler#1486 never-destroy rule for the buffer lifecycle, StagingStats route counters, the bit-identical staging invariant, and the before/after profile evidence owed on landing. Git integration base bumps to a456e6e. No product code in this change; the implementation commits follow the spec in the same pull request (recorded row shape). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki]
…persistent device buffer (mudler#2244) UploadRowsBf16 built a fresh ttnn::Tensor via from_span on every staging upload, so an identical-geometry staging paid tt-metal's full creation path each step: MeshBuffer::create on the device, MeshTensor/TensorAttributes construction, tensor-id assignment, and GraphTracker tracking, before the same bytes reached the same geometry. W4's profile attributed ~23% of the staging chain to that per-upload internal work. The bulk bf16 arm now stages through the slot's PERSISTENT device buffer. The first staging for a (slot, geometry) still runs the full from_span creation and the returned tensor is kept on the BufferSlot (which lives in the never-destroyed Slots() map, so the tensor is never destroyed after tt-metal teardown — mudler#1486). Every later staging of the same geometry packs the host bytes with the same function from_span calls (tt-metal host_tensor_from_span_with_pad_value, ttnn/core/tensor/tensor.cpp:170) and writes them through tt-metal's own in-place H2D — ttnn::copy_to_device into the resident MeshTensor (ttnn/core/tensor/tensor_ops.cpp:161), which takes MeshCommandQueue::enqueue_write/enqueue_write_shards against the existing buffer (tt_metal/impl/tensor/tensor_apis.cpp:149) instead of allocate_mesh_tensor_on_device_with_topology. The consumer-visible shadow in BufferSlot::device is unchanged: dropped by every host write, replaced by commits and reshapes; the resident buffer holds stale bytes that are only ever observed through a shadow a full staging write has just refreshed. Staging stays bit-identical: same packing function, same spec, same pad, a same-geometry buffer fully overwritten each time. The f32 arm keeps its declared dtype and never enters the route; interior views keep the anonymous from_span arm (W2c: never store against the base slot). The capture-unsafe host-write trace guards keep their semantics on both arms. StagingStats gains uploads_persistent_bf16, uploads_persistent_allocs and staged_persistent_bf16_bytes (uploads_bulk_bf16 keeps counting every bulk bf16 staging, so the W4 route pin is unchanged). The new case "kTENSTORRENT W5 EnsureDevice2D persistent staging buffer: route, reuse, bytes" pins cold allocation (2 allocs for 2 cold slots), reuse (a rewritten master restages with zero new allocations and the device readback carries the new bits), the geometry-change reallocation, and the f32 exclusion; it went red before the route existed (6 failing assertions, all zero counters) and seven mutations of the route — reuse disabled, write/alloc/byte counters dropped or halved, the device write skipped, the geometry guard removed, the slot store dropped — each fail the focused case and were restored byte-for-byte. Full test_tenstorrent_backend: 45/45 cases, 5062/5062 assertions on the P150 under the gpu.lock. The W4 evidence owed by the spec (before/after profile on the P150 and the benchmark-record entry) stays with the row; this wave changes speed, never tokens. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki]
… and the trace moves the wall's address (mudler#2244) The wave's own gate said what landing means here: the capability is real and the speed claim is not. Both are recorded, because "a wall that does not move is a reported result, not a failure" is the spec's term for exactly this outcome. Operator evidence on dc473a94c, one lock hold per series: - Correctness: e2e leg A eager 16/16 STRICT token-exact, leg B ambient 16/16 STRICT token-exact, suite 45/45 cases / 5062 assertions exit 0 — the sacred pair is byte-identical; preflight 0 FAIL with the three standing asset SKIPs. - Review: fresh reviewer PASS (8 mutations including the reachability mutation; full gate rerun on the immutable head). One LOW comment-only finding rides the next change to tenstorrent_device.h; one W6+ aliasing awareness note is recorded in the evidence log. - Speed: interleaved same-method A/B (fresh scratch build of a456e6e vs dc473a94c, JIT-discard per arm, perf per measured leg) — 19.154 s vs 19.181 s for 3 tokens, -0.14%, noise. The profile then split the W4 hypothesis: allocate_mesh_tensor is 0.02% of the AFTER profile and the enqueue_write stacks are identical in both arms, so the per-upload allocation was never the wall; the wall is the per-CQ-operation tt-metal stack (MetalContext::instance 11.14%, Cluster::get_chip 5.90%, read_cq_host_ptr 5.27% plus sub-slices) and Threadpool::PollForWork 14.29%. Record moves: ## Now records W5's landing and the honest unmoved wall; ## Evidence gains the W5 entry with the A/B and the attribution; the owed next lever becomes W6 (batch per-layer staging — one CQ write per step divides the per-op tax by the fan-in, our file set), with the tt-metal-side residual recorded as the upstream-shaped alternative. The full log is docs/bench-evidence/tt-qwen35-eager-profile-w5-20260829.log (forced add; evidence logs are gitignored). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:zai-glm-5.3-flash [maki]
…head_count_kv`, and stop counting the MTP block as a layer (mudler#2243, mudler#2177) (mudler#2269) fix(MODEL-MM-GLM53-FLASH): read the layer schedule out of `attention.head_count_kv`, and stop counting the MTP block as a layer `Glm5NextHfConfigFromGguf` read `glm5next.attention.head_count_kv` as a scalar and then REQUIRED a `glm5next.layer_types` string array that only `scripts/convert-glm5-next-gguf.py` writes. The published `unsloth/GLM-5.3-Flash-GGUF` UD-Q2_K_XL artifact carries neither shape, so the production loader opened all four shards, sized all 1412 tensors and stopped on a config key with `key glm5next.attention.head_count_kv is not an integer`. ## Two issues, one defect mudler#2243 and mudler#2177 are the same fault seen from two sides, and fixing either alone leaves the other standing. The crash is the array spelling. The deeper fault is that the schedule was never read at all, and the `idx % 4 != 3` pattern that would have stood in for it is right on this checkpoint by coincidence, so a fine-tune that moves one layer would build a fluent wrong model with no gate able to see it. They are therefore fixed together. ## What the artifact actually stores Read out of shard 1's own KV block, all 72 keys, 2026-08-29: ```text glm5next.block_count : u32 = 46 glm5next.nextn_predict_layers : u32 = 1 glm5next.attention.head_count : u32 = 64 glm5next.attention.head_count_kv : array[i32] n=46 glm5next.swiglu_clamp_exp : array[f32] n=46, every entry 10.0 glm5next.swiglu_clamp_shexp : array[f32] n=46, every entry 10.0 glm5next.layer_types : ABSENT ``` The `head_count_kv` array is **34 zeros and 12 ones**, the ones at 3, 7, ..., 43 **and 45**. ## Reading the schedule The builder accepts llama.cpp's scalar-or-array spelling of the key — `b10451:src/llama-model.cpp:1177` reads it through `get_key_or_arr(LLM_KV_ATTENTION_HEAD_COUNT_KV, hparams.n_head_kv_arr, hparams.n_layer(), false)` — and DERIVES the schedule from the values with llama.cpp's own predicate, `is_recr_impl[i] = hparams.n_head_kv(i) == 0`, spelled for this model's KDA parent at `b10451:src/models/kimi-linear.cpp:18` with the comment "KDA layers are recurrent". `swiglu_clamp_exp` and `swiglu_clamp_shexp` are read in the same scalar-or-array shape, and `swiglu_clamp_shexp` is read at all for the first time. **The array is a schedule and not a KV-head count.** Its non-zero entries are `1`, the single latent KV head MLA has, while upstream requires `num_attention_heads == num_key_value_heads` and the released `config.json` states 64 for both. Assigning `1` would refuse the published file with a true statement about a number that file never makes, so the array form leaves `num_key_value_heads` at upstream's `None -> num_attention_heads` default. Four refusals rather than a tolerant reader, because each is a shape that would otherwise load quietly: a `layer_types` that disagrees with `head_count_kv` refuses at the first block where they differ, compared on the layer KIND so upstream's `full_attention` spelling is not a false positive; a per-block array whose length is not `block_count` refuses with the key and the shape found; a non-uniform clamp array refuses, because upstream has ONE `swiglu_limit`; and a file that states the schedule in neither spelling refuses, naming both keys. That last one is where the removed `ReqStrArray`'s obligation went — it moved, it was not weakened. ## Blocks are not layers Review caught the second half. `c.num_hidden_layers` was set straight from `block_count`, so ONE model resolved to a 45-layer backbone from its `config.json` and a 46-layer one from its GGUF. The extra entry is the multi-token-prediction block, and the first draft of this PR pinned the 46 as correct — the shape where a gate stops being able to see a defect. Nothing downstream would have refused it. `ParseGlm5NextParams` sizes all three schedules from `num_hidden_layers`, so W5b (mudler#2241) and W5c (mudler#2242) would have built a decoder layer out of the MTP block. It would run and produce plausible tokens. The contract is BACKBONE depth, and the tree already said so twice: `glm5_next.h:193` annotates the field `// 45`, and the existing inventory case asserts `CHECK_FALSE(has("blk.45.attn_norm.weight"))` because the reference discards the MTP block. llama.cpp states the relationship in its own converters, at the pinned RELEASE and not only in the `glm5next` pull request: ```python self.block_count = self.hparams["num_hidden_layers"] + self.hparams.get("num_nextn_predict_layers", 0) # b10451:conversion/exaone.py:134, and the same += at conversion/deepseek.py:470 and :545 ``` So `num_hidden_layers = block_count - nextn_predict_layers`; every per-block array is validated against `block_count`, which is the length llama.cpp's own `get_key_or_arr(..., n_layer, false)` reads; the three schedules are truncated to the backbone; and a file claiming more MTP blocks than blocks is refused by name. Our converter writes `block_count = n_layers` with `nextn_predict_layers = 0`, so its output is unchanged. The MTP block is read, counted and DROPPED. No field on `HfConfig` or `Glm5NextParams` carries `nextn_predict_layers`, deliberately — the head is O2's and W5b's — and the spec records what W5b inherits. ## The assertions `CHECK(p.num_hidden_layers == 46)` is gone. In its place: ```cpp constexpr int64_t kBlocks = 46; constexpr int64_t kMtp = 1; CHECK(p.num_hidden_layers == kBlocks - kMtp); CHECK(p.num_hidden_layers == 45); // and 45 is what config.json declares CHECK(static_cast<int64_t>(p.layer_types.size()) == kBlocks - kMtp); CHECK(static_cast<int64_t>(p.mlp_layer_types.size()) == kBlocks - kMtp); CHECK(static_cast<int64_t>(p.indexer_types.size()) == kBlocks - kMtp); CHECK(p.layer_types == KindsOf(backbone)); // the array's first 45 entries CHECK(kv[44] == 0); CHECK(kv[45] == 1); CHECK(p.layer_types.back() == Glm5NextLayerKind::kLinearAttention); CHECK(p.num_dsa_layers() == 11); // 12 MLA-shaped BLOCKS, 11 layers ``` `layer_types.back()` is the sharp one: entry 44 is a `0` and entry 45 a `1`, so a reader that forgets to truncate ends its stack with a DSA layer built out of the MTP block. The same 46-entry file with `nextn_predict_layers = 0` resolves to 46 layers and twelve MLA-shaped ones, which is what makes the subtraction observable rather than decorative. And a new case asserts the cross-source agreement directly: a GGUF and the released `config.json` of this one model now resolve to the same depth, the same schedule and the same KDA/DSA counts. ## Evidence **Red 1 — the array keys**, all six of the first cases failing before the parse change, the two headline ones with the production message the issue quotes: ```text ERROR: test case THREW exception: glm5_next gguf: key glm5next.attention.head_count_kv is not an integer [doctest] test cases: 24 | 18 passed | 6 failed | 0 skipped [doctest] assertions: 2353 | 2343 passed | 10 failed | ``` **Red 2 — the other reading of `block_count`.** With `n_layers = n_blocks` and everything else identical, 15 assertions fail across 3 cases: ```text CHECK( p.num_hidden_layers == kBlocks - kMtp ) NOT correct! CHECK( p.layer_types.back() == Glm5NextLayerKind::kLinearAttention ) NOT correct! CHECK( p.num_dsa_layers() == 11 ) NOT correct! CHECK( from_gguf.num_hidden_layers == from_json.num_hidden_layers ) NOT correct! CHECK( from_gguf.layer_types == from_json.layer_types ) NOT correct! [doctest] test cases: 26 | 23 passed | 3 failed | 0 skipped [doctest] assertions: 2403 | 2388 passed | 15 failed | ``` **Green**: `26 | 26 passed | 0 failed | 0 skipped`, `assertions: 2403 | 2403 passed | 0 failed`. **Ten mutations, each built and run, each detected**, tree restored byte-for-byte: | Mutation | Detected by | |---|---| | derive -> synthesize `idx % 4 != 3` | the non-stride file proves it | | delete the `layer_types` / `head_count_kv` cross-check | `layer_types` and `head_count_kv` must agree | | delete the `head_count_kv` length check | a per-layer array whose length is not `block_count` | | accept a non-uniform clamp array | `swiglu_clamp_exp` and `_shexp` in the ARRAY form | | ignore `swiglu_clamp_shexp` | `swiglu_clamp_exp` and `_shexp` in the ARRAY form | | default the schedule instead of refusing | a GGUF that states NO schedule | | `block_count` straight into `num_hidden_layers` | a GGUF and a config.json resolve identically | | keep the MTP block as a 46th layer | the published GGUF states its schedule ONLY in `head_count_kv` | | validate per-block arrays against the backbone depth | the published GGUF states its schedule ONLY in `head_count_kv` | | delete the more-MTP-than-blocks refusal | more MTP blocks than blocks is refused | **The non-stride fixture is what proves the values are read.** Over the 45 backbone layers the published array agrees with `idx % 4 == 3` exactly — the case asserts that agreement rather than leaving it implied — so a fixture whose array is `[1, 0, 1, 0, 0, 0, 0, 1]`, disagreeing on three of eight blocks, is the one a synthesized schedule cannot pass. ## The loader's new stopping point Measured on the staged artifact through `LoadedEngine::FromModelDir` on `device = kCPU`, headers only, with the change reverted and restored so the before/after is ONE binary and one tree, and re-measured unchanged after the `block_count` fix: ```text without the fix : glm5_next gguf: key glm5next.attention.head_count_kv is not an integer with the fix : vt: glm5_next gguf: attention.key_length_mla - attention.key_length is -256 but rope.dimension_count is 0; the file states this model's rotary width twice and the two disagree ``` That file is not malformed. llama.cpp writes `key_length = kv_lora_rank + qk_rope_head_dim` and `key_length_mla = qk_nope_head_dim + qk_rope_head_dim` (`b10451:conversion/deepseek.py:345-348`), which for this model gives the artifact's own 512 and 256; our converter writes `key_length = qk_nope_head_dim`, a different quantity under the same name. `glm5next.attention.linear_head_count`, a `ReqInt` in the builder, is in none of the file's 72 keys and llama.cpp spells it nowhere. Both move the WRITE side, so they are filed as mudler#2268 and owned by this row. That is the next milestone. ## Records `.agents/specs/glm5-next-flash.md` O18 records the three arrays as discharged, carries the corrected 34 / 12 block structure, states the block-versus-layer contract and what W5b inherits from it, and names the new stopping point and its cause. Two append-only rows in `.agents/issue-index.md`, for mudler#2177 and for mudler#2268; mudler#2243 already had one. No existing row or O-number was edited or renumbered. The mudler#2268 row quotes a `file:line` anchor that moved later in this PR; O18 is the corrected surface and says so. ## Gates `scripts/agent-preflight.sh --fail-on-skip`: **109 gates ok, 0 skipped, 0 failed, "All gates green."** FOLLOWING_AGENTS_PROTOCOL Closes mudler#2243 Closes mudler#2177 Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…pport Add PagedAttnDecodeGqaF32Q — a warp-strided decode attention kernel for the f32-query x bf16/fp8-KV x f32-output dtype combination that the GGUF dense path uses. Without this arm, f32-query decode falls through to PagedAttnOnline, which walks each context token with a per-key __syncthreads() reduction (~593us/call, 8 calls/token on Qwen3.5-4B). The kernel fuses QG=4 query heads per KV group, walks the sequence warp-strided, and does register online softmax with no per-key block sync. EPL = d/32 (8 at d=256, 4 at d=128). Dispatch is gated behind VT_ATTN_DECODE_GQA4=1 for the exact geometry (f32 q/out, bf16 or fp8 k/v, d=128 or 256, hq=16, kv=4). fp8 KV cache support: template PagedAttnDecodeGqaF32Q on TKV (default __hip_bfloat16) with k_scale/v_scale params. Add LoadRowEplFp8<EPL> (vectorized uint8_t loads + F8E4M3ToF32Dev dequant x scale) and LoadRowEplKv<EPL, TKV> (if constexpr dispatch between bf16 and fp8). Widen the dispatch guard: kBF16 || (kI8 && kv_cache_dtype != kAuto). The dequant arithmetic is identical to LoadKv(uint8_t*, ...) at line 176 — the fast kernel calls the same dequant the fallback already used. Benchmark: Qwen3.5-4B Q4_K_M, RX 7900 XTX, ROCm 7.14.0, 128-token greedy decode, 4 reps: ctx=256: 99.94 -> 140.19 tok/s (1.40x), bf16=143.0 (0.98x) ctx=1024: 56.28 -> 122.31 tok/s (2.17x), bf16=128.7 (0.95x) ctx=4096: 20.53 -> 79.53 tok/s (3.87x), bf16=92.7 (0.86x) ctx=8192: 11.08 -> 54.55 tok/s (4.93x), bf16=66.4 (0.82x) ctx=16384: 5.78 -> 33.47 tok/s (5.79x), bf16=43.16 (0.78x) Token-exact: 3/3 simple prompts match the PagedAttnOnline fallback output byte-for-byte. Divergence at thinking-tag near-ties is the reduction-order risk the bf16 GQA4 path already carries — same policy class as VT_ATTN_DECODE_D128. The arm ships opt-in (VT_ATTN_DECODE_GQA4=1); near-tie adjudication before any default-ON flip remains owed. Spec: .agents/specs/rocm-fp8-kv-decode-attn.md Evidence: docs/bench-evidence/gfx1100-tg200-t3a-20260823.md Fork issue: #7 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP]
W6 store/read is on main. This branch's PagedAttnDecodeGqaF32Q arm reads fp8 KV. Without advertising the dtypes, --kv-cache-dtype fp8 is refused at the backend registry before the kernel can run. Fork issue: #7 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-3-FLASH [OMP]
…anchor Integrating upstream main at c3522bc brought the record-anchor ratchet onto this branch, and it flagged two regressions this PR owns. The GFX1100-TG200 row in `.agents/issue-index.md` quoted the widened `VT_ATTN_DECODE_GQA4` guard as `kBF16 || kI8` with the pipes unescaped inside the code span, so the table-shape gate counted 7 pipes against the table's 5 and the row mis-rendered on GitHub (mudler#1033's exact shape). The pipes are now `\|\|`, which the cell splitter and the pipe counter both treat as escaped. The `supported_kv_cache_dtypes` insertion in `include/vllm/v1/attention/backend.h` shifted `get_kv_cache_shape` from :580 to :587, staling the `KERNEL-ATTN-MLA-SPARSE` citation by exactly one over the ratchet's baseline (29 stale against 28). The citation is advanced to :587; the baseline is untouched. `scripts/check-agent-record.py` exits 0: ENGINE=173 MODEL=379 QUANT=85 KERNEL=58 BACKEND=87 ANCHOR-ROT=33. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-3-FLASH [OMP]
`check-env-doc` reads `std::getenv("VT_ATTN_DECODE_GQA4")` in
`src/vt/rocm/rocm_paged_attn.hip` and the knob was on neither surface:
not documented in `docs/ENVIRONMENT.md` and not on the kernel-internal
allowlist. It is a behavior-changing dispatch gate, not a tuning
switch, so the row documents what setting it routes where, the exact
geometry it guards, and the near-tie caveat the read site's own comment
carries.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:GLM-5-3-FLASH [OMP]
ghazni101
force-pushed
the
row/fp8-kv-decode-attn
branch
from
August 29, 2026 13:28
74510c5 to
a05d760
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What changed
Add
PagedAttnDecodeGqaF32Q— a warp-strided decode attention kernel forthe f32-query x bf16/fp8-KV x f32-output dtype combination that the GGUF
dense path uses. Without this arm, f32-query decode falls through to
PagedAttnOnline, which walks each context token with a per-key__syncthreads()reduction (~593us/call, 8 calls/token on Qwen3.5-4B).The kernel fuses QG=4 query heads per KV group, walks the sequence
warp-strided, and does register online softmax with no per-key block
sync. EPL = d/32 (8 at d=256, 4 at d=128). Dispatch is gated behind
VT_ATTN_DECODE_GQA4=1for the exact geometry (f32 q/out, bf16 or fp8k/v, d=128 or 256, hq=16, kv=4).
fp8 KV cache support
--kv-cache-dtype fp8set the KV cache toDType::kI8, which failedthe
kBF16guard on every optimized decode kernel inrocm_paged_attn.hip. The dispatch fell through toPagedAttnOnline— the reference kernel with per-key
__syncthreadsreduction, no GQAamortization, no vectorized loads. Up to 7.5x decode regression at long
context on gfx1100.
Fix: template
PagedAttnDecodeGqaF32QonTKV(default__hip_bfloat16) withk_scale/v_scaleparams. AddLoadRowEplFp8<EPL>(vectorized uint8_t loads +F8E4M3ToF32Devdequant x scale) and
LoadRowEplKv<EPL, TKV>(if constexprdispatchbetween bf16 and fp8). Widen the dispatch guard:
kBF16 || (kI8 && kv_cache_dtype != kAuto).The dequant arithmetic is identical to the existing
LoadKv(uint8_t *, ...)at line 176 — the fast kernel now calls the same dequant thefallback already used.
Benchmark
Qwen3.5-4B Q4_K_M, RX 7900 XTX, ROCm 7.14.0, 128-token greedy decode,
4 reps:
Correctness
3/3 simple prompts produce token-exact identical output vs the
PagedAttnOnlinefallback. Divergence at thinking-tag near-ties is thereduction-order risk the bf16 GQA4 path already carries — same policy
class as
VT_ATTN_DECODE_D128. The arm ships opt-in(
VT_ATTN_DECODE_GQA4=1); near-tie adjudication before any default-ONflip remains owed.
Relationship with mudler#2080 and mudler#1936
This PR is the fp8 decode-attention split from mudler#1936, per the reviewer's
request. mudler#1936 carried two units of work: the GFX1100-TG200 decode
campaign and this fp8 KV decode-attention arm. This PR contains only the
fp8 arm.
This PR builds on the ROCm fp8-e4m3 KV cache store and read that landed
in mudler#2080 (
191f64608). The store side is onmain; this PR adds thedecode-attention read side. The
08595f335fix (mudler#2161) for the G1btest predicate is also in the base.
The TG200 campaign remains on mudler#1936 and will be sent separately after
this lands.
How to verify
cmake -B /tmp/build -DVLLM_CPP_HIP=ON make -j4 vllm test_rocm_fp8_kv_cache ctest -R 'rocm_fp8'Verified on gfx1100 (RX 7900 XTX), ROCm 7.14.0:
test_rocm_fp8_kv_cache: 7/7 cases, 28 assertionstest_rocm_backend: 9/9 cases, 1065 assertionstest_ops_quant_dot: 28/28 cases, 210138 assertionsOut of scope
fp8_e5m2 compute, per-attention-head scales, the Metal arm, fast-path
(tensor-core/rocWMMA) fp8 attention kernels, the bf16 decode-opt
kernels, the prefill fp8 path, non-gfx1100 architectures.
Spec:
.agents/specs/rocm-fp8-kv-decode-attn.mdEvidence:
docs/bench-evidence/gfx1100-tg200-t3a-20260823.mdTracking issue: #7
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:GLM-5-2 [OMP]