perf(gemma4): dual-GPU FP8 resident without host OOM + peer mix - #154
Conversation
Warm decode: MoE ~95% of layer time (attn~4% mlp~1%). Microbench: GemmEx BF16 gate_up~30us down~17us; hipBLASLt FP8 heuristic unsupported on gfx1201. Upstream check: mudler#154 open; decode-graph is CUDA-only (ROCm no graph capture). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
…re-push sandbox (#159) Two guards on main were RED and between them blocked every open external contributor PR (#127, #154, #155) and every push. Both premises were verified in the tree before changing anything. 1) check-device-leakage: src/vllm/v1/worker/gpu/runner.cpp named vt::DeviceType::kCUDA in the device-agnostic shared layer (DSR bucket 'kcuda' 1 > baseline 0). It came in with the QueueSupportsAsyncInputCombine rescope during the PR #140 fix round - ours, not the contributors'; richiejp reported it in #127's honest gaps. Fixed the way the guard's own message prescribes, mirroring the SupportsAuxStream precedent: ask the backend, not the device. New vt::Backend::SupportsAsyncSampledTokenReadback() (base false) answers whether the host may validly read the sampled token id back between steps; CPU overrides true (host and device memory are one allocation) and CUDA overrides true (the id is device-mirrored). The runner asks vt::TryGetBackend(queue.device.type), whose nullptr for a device absent from the build also subsumes the old #ifdef VLLM_CPP_CUDA guard. SEMANTICS UNCHANGED: CPU async-ON, CUDA async-ON, discrete non-CUDA (ROCm gfx1201) async-OFF - the "!"-token hazard stays closed. 2) .githooks/pre-push ran check-policy.py inside a PARTIAL export (README.md docs scripts .agents), but policy_contract.py:428 asserts AGENTS.md is a non-symlink regular file and resolves its Markdown links against that sandbox. AGENTS.md and its .env.example link were both missing, so the hook failed closed on content that is fine in the real tree - every push refused. EXPORT_PATHS is now a superset of what the CHECKERS read. Gates: check-device-leakage RED->GREEN (kcuda=0, DSR 32 == baseline 32); all four hook checkers OK in the reproduced sandbox; test_async_llm 8/8-347, test_engine_core 6/6-44, test_llm_engine 11/11-204 (CPU still resolves async-ON); clean -Werror CPU build; full 11-gate record battery green. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Claude (Opus 5) via Claude Code
localai-bot
left a comment
There was a problem hiding this comment.
Thanks for the follow-up to #140 — the dual-GPU FP8 resident direction is right, and the peer-mix idea is one we want. Two things to sort before this can land, one mechanical and one structural. They have the same fix.
1. The CPU, CUDA and Vulkan builds do not link
build-test-cpu and build-test-vulkan are failing for a real reason, not a flake:
gemma4_moe.cpp: undefined reference to `vllm::PinGemma4Fp8ExpertHostCache(...)`
gemma4_moe.cpp: undefined reference to `vllm::PeerCopyGemma4ExpertSlice(...)`
gemma4_moe.cpp: undefined reference to `vt::rocm::MatmulBTFp8ChannelRocm(...)`
gemma4_moe.cpp: undefined reference to `vt::rocm::ExpertGeGLUBf16TopKM1Rocm(...)`
gemma4.cpp: undefined reference to `vt::rocm::RmsNormPlusAddRocm(...)`
gemma4.cpp: undefined reference to `vt::rocm::DualRmsNormPlusResRocm(...)`
gemma4.cpp: undefined reference to `vt::rocm::GeluMulSeparateRocm(...)`
These are declared in include/vllm/model_executor/models/gemma4_moe.h but defined only in src/vt/rocm/rocm_gemma4_experts.hip, and that file is added to the build only under the ROCm branch (CMakeLists.txt:1148). Every non-ROCm configuration compiles the calls and then has nothing to link them against.
2. Backend names in the device-agnostic layer
The deeper reason it breaks is that src/vllm/model_executor/models/ is the shared, device-agnostic layer, and this PR puts 11 vt::rocm:: call sites into it (4 in gemma4.cpp, 7 in gemma4_moe.cpp). Today main has exactly zero — the only vt::rocm:: reference outside src/vt/ is in src/vllm/platforms/rocm.cpp, which is the right place for it. This is what check-device-leakage guards, and it is a rule we hold ourselves to: I had to fix one of my own violations in #159 this week for the same reason.
(Do rerun CI after rebasing, by the way — the device-leakage failure you are seeing right now is partly our bug, not yours. It reports kcuda=1 from a DeviceType::kCUDA I left in runner.cpp. That is fixed on main in #159.)
The fix for both
Route these through the ops/backend seam rather than calling the ROCm entry points by name:
- Kernels (
RmsNormPlusAddRocm,DualRmsNormPlusResRocm,GeluMulSeparateRocm,MatmulBTAlphaBetaRocm,MatmulBTFp8ChannelRocm,ExpertGeGLUBf16TopKM1Rocm): register them asOpIdimplementations forDeviceType::kROCmand call the portablevt::entry point. The model file then reads the same on every backend, and CPU keeps its reference implementation. - Policy/capability (
PeerCopyGemma4ExpertSlice,PinGemma4Fp8ExpertHostCache): these are "can this backend do peer copy / pinned expert staging", so they belong onvt::Backendas virtuals with a conservative base implementation, in the shape ofSupportsAuxStream(include/vt/backend.h:122) andSupportsAsyncSampledTokenReadback(:141). Non-ROCm backends inherit the base and the peer path simply does not engage.
That keeps your ROCm fast path exactly as fast, and makes every other build link again.
Rebase onto current main while you are in there — the merge base here is c05cee1d, several commits back — and agent-record / documentation-checkpoint should be satisfied per AGENTS.md (a record under .agents/ bound to this row, plus the docs/ checkpoint files).
Happy to help with the ops-registration wiring if it would be useful — say the word.
|
Thanks for the clear CHANGES_REQUESTED review — fully agree on both points. Ack
Process / PR hygiene
Rebasing/reworking the ROCm registration path next. Happy to take a pointer if there's a preferred OpId naming pattern for the fused RmsNorm/GeGLU helpers beyond matching existing MatmulBT-style registrations. FOLLOWING_AGENTS_PROTOCOL |
Address mudler#154 review: device-agnostic model layer no longer names ROCm. Fused helpers dispatch ROCm kernels under VLLM_CPP_HIP; non-HIP peer/pin stubs. check-device-leakage holds baseline. Docs STATUS/BENCHMARKS/FEATURES/USAGE. Drop PR_DRAFT_GEMMA4_RESIDENT_SPEED.md from tree (pr-size unclassified path). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Address mudler#154 review: device-agnostic model layer no longer names ROCm. Fused helpers in include/vt/fused_ops.h (ROCm under VLLM_CPP_HIP; non-HIP stubs). Peer/pin stubs off-HIP. check-device-leakage holds baseline. Drop PR_DRAFT_GEMMA4_RESIDENT_SPEED.md (pr-size). Docs STATUS/BENCHMARKS/FEATURES/USAGE/README. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
e6a314b to
324c405
Compare
Progress per layer, 3GiB headroom on compute GPU, fill GPU0 before GPU1 for same-device fast path. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Resident path dequanted every expert into permanent host caches (~1.5GiB/layer), OOM on ~30G RAM. Stream per-expert ephemeral dequant + H2D; clear caches. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Stream per-expert dequant (no permanent host BF16 cache). 10GiB headroom on compute GPU. Off-device resident layers use ephemeral H2D instead of double-alloc. Lab: 30 layers 42.5GiB resident, Paris stop EXIT=0 peak VRAM ~28.6GB. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Enable HIP peer access; GPU1 expert slices peer/stage to compute scratch instead of FP8 host dequant. Lab: full 30L resident Paris EXIT=0, wall ~198s (was ~294s). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Default on-GPU weighted accumulate; VT_GEMMA4_HOST_AXPY=1 falls back. Lab: full resident Paris EXIT=0 ~191s wall. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Allocate peer/weight/yscale scratch once per token instead of per expert. Reduces HIP alloc churn during decode. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Stream path no longer reserves ~12MiB expert peer buffers per layer step. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
After one model load, run N blocking completions with per-run secs and tok_s. Lab (Firworks FP8 stream, gfx1201): run1 ~0.14 tok/s (first-use expert dequant), run2 ~24 tok/s warm. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Row-parallel DequantFp8ChannelToBf16; parallel top-k expert cache fill when cold; nesting guard avoids thread storms. Lab: cold ~0.3 tok/s (was ~0.14), warm ~4 tok/s, Paris OK. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
hipHostRegister after first dequant for faster H2D fallback; log device expert OOM fallbacks. Lab: warm still ~3.8-4 tok/s with successful device expert path (Paris OK). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Shared gate/up/gu/act/weight buffers per token; T=1 pack is two bulk copies. Lab: warm still ~3.8 tok/s (GEMM-bound); Paris OK. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
hipblasGemmStridedBatchedEx + GemmBatchedEx pointer path for top-k expert fuse. Gemma MoE uses VT_GEMMA4_BATCH_EXPERTS=1 to enable; default remains serial ExpertGeGLU (batched path currently slower ~0.7 vs ~3 tok/s warm on R9700 due to Gelu pack). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Single GeluAndMul over [G,2I] + pointer-batch down AB. Default remains serial (VT_GEMMA4_BATCH_EXPERTS=1 to try). Lab: fused-batch warm ~0.5 tok/s vs serial ~3. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
ExpertScratch, peer/tmp buffers, y/ysum live for the layer call so multi-token prefill and decode avoid re-alloc. Lab warm ~4.5 tok/s (was ~3.9). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Ungrouped softmax+greedy top-k HIP kernel; Gemma MoE only D2Hs [T,K] weights/indices (not full [T,E] logits). Lab: Paris OK, warm ~4.4 tok/s. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Lab stream FP8 --repeat2: warm mean ~7.6 tok/s (paris 6.8, arith 11), 3/3 correct. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
- GeluMulSeparateRocm: no gate|up pack before GeGLU (ExpertGeGLU hot path) - thread_local MoE scratch reused across 30 layers (pool thrash fix) Lab stream FP8 --repeat2: 3/3 OK, warm mean ~18.9 tok/s (paris 26.7). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
MatmulBTAlphaBetaRocm: ysum = w*(act@W^T) + beta*ysum — drops MulScalar+Add on device path. Lab: 3/3 OK, warm mean ~18.7 tok/s (paris 24.8). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Lab: 3/3 OK, warm mean ~19.7 tok/s (paris 27.0). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Reuse dhn/h1/h2/moe/n*/PLE temps across 30 layers; residual stays on hidden via D2D publish. Lab: 3/3 OK, warm mean ~18.8 tok/s (paris 25.2). FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Skip vt::MatmulBT ops dispatch on gate_up loop. Lab: Paris OK, warm ~35.2-35.4 tok/s. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
…M_EXPERT) Multi-block act (gelu*up) + down accumulate. Wired into MoE top-k path. Lab: Paris OK; warm ~34.2 vs hipBLAS ~35.2. Default remains hipBLAS. rocWMMA probe OK on gfx1201 for future MFMA pass. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
- LoadChatTemplateFromConfig falls back to sibling chat_template.jinja (Gemma4 HF layout) - VT_SERVER_VERBOSE=1 / --verbose: log roles, prompt preview, token counts, output preview FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
api POST size, templated, stream/generate begin, per-second step progress, done. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
- RequestLogger: Received/Finished/api ingress/errors (enable-log-requests) - --enable-log-outputs --max-log-len --enable-metrics - Wire PrometheusStatLogger to GET /metrics - Keep chat-dbg stages under --verbose - Default log-requests + metrics ON for solid agent harness FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Default VT_SERVER_MAX_PROMPT_CHARS=48000, VT_SERVER_MAX_NEW_TOKENS=4096. Prevents async prefill wedge on 140k-char full-agent system dumps. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Pairs with start-gemma4-fp8-8010.sh max_model_len=262144 + prefix cache. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Enabled via VT_SERVER_PREFILL_PROGRESS=1 or VT_SERVER_VERBOSE=1; ~2Hz rate limit. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Root cause of chat hang under concurrent Hermes: LLMEngine::generate looped while has_unfinished_requests() globally, so one huge async job blocked all blocking generate forever. Wait only until this request_id finishes. Chat serving (Gemma4/Hermes, vLLM recipe parity): - apply_chat_template / MakeChatTemplatePromptFn take enable_thinking - --enable-thinking / --no-enable-thinking (default off for agents) - empty thought block when thinking off (HF jinja) Lab: chat count max=48 finishes; warm chat ~32 tok/s after expert cache. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
Address mudler#154 review: device-agnostic model layer no longer names ROCm. Fused helpers in include/vt/fused_ops.h (ROCm under VLLM_CPP_HIP; non-HIP stubs). Peer/pin stubs off-HIP. check-device-leakage holds baseline. Drop PR_DRAFT_GEMMA4_RESIDENT_SPEED.md (pr-size). Docs STATUS/BENCHMARKS/FEATURES/USAGE/README. FOLLOWING_AGENTS_PROTOCOL Assisted-by: Hermes:grok-4.5 [Hermes]
324c405 to
a6e1ecf
Compare
Address mudler#154 review: device-agnostic model layer no longer names ROCm. Fused helpers in include/vt/fused_ops.h (ROCm under VLLM_CPP_HIP; non-HIP stubs). Peer/pin stubs off-HIP. check-device-leakage holds baseline. Drop PR_DRAFT_GEMMA4_RESIDENT_SPEED.md (pr-size). Docs STATUS/BENCHMARKS/FEATURES/USAGE/README. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.5 [Hermes]
a6e1ecf to
e086435
Compare
Progress update (post-review)Addressed the structural review items we could land without a full PR split:
Still red (honest)
Split already doneGeneric serve hang fix is #163 (not blocked on this MoE campaign). Next for #154 when size policy is the priority: squash or carve into ≤900-line PRs (ops seam first, then expert path, then kernels). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true |
localai-bot
left a comment
There was a problem hiding this comment.
Re-reviewed at e086435e — the defect I flagged is fixed, and fixed the right way.
Verified rather than taken on trust:
git grep -c "vt::rocm::" -- src/vllm/model_executor/on your head returns zero (was 11 acrossgemma4.cppandgemma4_moe.cpp).- The new
include/vt/fused_ops.h+src/vt/fused_ops.cppseam carries exactly the right contract in its own header comment: "Model code MUST call these — nevervt::rocm::*— so CPU/CUDA/Vulkan link." Non-HIP stubs sit under the#endif // VLLM_CPP_HIPguard. - CI on that SHA:
build-test-cpu✅,build-test-vulkan✅,device-leakage✅, both sanitizers ✅. The three link failures I quoted are gone.
That is the ops/backend-seam route rather than a #ifdef around the call sites, which is what keeps the device-agnostic layer honest. Nice.
Remaining red is process only — pr-size, agent-record, documentation-checkpoint. Per operator instruction I am merging and carrying those myself rather than bouncing the PR back to you again.
Thanks for the ROCm campaign generally — the gfx1201 evidence in #41 is the strongest data point we have on the discrete lane.
check-public-doc-tables and check-env-doc are red on pristine origin/main and block every push through the pre-push hook. Neither came from a feature branch; both arrived with #154. This fixes the two conditions I can judge correctly and deliberately leaves the rest rather than guessing at another row's intent. FIXED. docs/STATUS.md carried 12 h2 sections against an 11 ratchet. The 12th was a DATED per-change narrative -- exactly what the checker says to collapse -- so its binding result becomes one line under the existing "Backend detail" section, which already owns per-backend state. Nothing is lost: the full narrative stays in the append-only record and the PR. Sections 12 -> 11, STATUS 277416 -> 277217 chars. VT_SERVER_VERBOSE is documented rather than allowlisted because it is user-facing: examples/server sets it from its own verbosity flag, and it is the umbrella switch VT_SERVER_PREFILL_PROGRESS falls back to when unset. NOT FIXED, and named so the next session does not rediscover them: * STATUS is still 277217 against a 276960 ratchet -- main was 456 over before this change and is 257 over after it. Closing that means cutting content I do not own, and the ratchet may only shrink, so it needs the owning row. * VT_GEMMA4_BATCH_EXPERTS, VT_GEMMA4_CUSTOM_EXPERT and VT_GEMMA4_EXPERT_VRAM_MB are undocumented. Each is either a user-facing knob for docs/ENVIRONMENT.md or a kernel-internal switch for the allowlist, and that is a statement about #154's intent. Guessing it would put a wrong claim on a public page, which is worse than leaving the gate red. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
docs/BENCHMARKS.md is a KEYED TABLE: a checkpoint updates a ROW, it does not append a dated H2 section. #154 appended one, which left main red on check-public-doc-tables for two reasons at once -- the non-canonical section, and the em-dash it carried (house style forbids em-dashes on the public pages). scripts/roll-benchmark-record.py --apply moves it verbatim into .agents/benchmark-record.md, which is where per-change narrative belongs, and both errors go with it. Nothing is lost and nothing is rewritten: the section moves byte-for-byte into the append-only record. This is main's red, not this branch's, and it was blocking the agent-record CI job on every open PR rather than just on the change that introduced it. RESIDUE, NAMED RATHER THAN PAPERED OVER. check-public-doc-tables still fails on one item: docs/STATUS.md is 277213 chars against a 276960 ratchet. This branch SHRINKS that page by 4 chars; main is already 257 over. Clearing it means collapsing superseded narrative, and the only block big enough to matter is a single 33,211-char table cell (the Laguna-S-2.1 MoE row) that is itself well past the 220-char cell rule. Collapsing a cell that size is a deliberate, separately-reviewable change with its own owner, not something to bury in a Vulkan performance PR, so it is left open and stated here instead. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
perf(vulkan): fused attn preamble native -- 16 host round trips/token gone kAttnQkNormRopeGate gets a native Vulkan kernel, removing the last reference-tier op that fires during 27B decode. Module count 24 -> 25; the Vulkan decline list is now exactly kRopeCosSinCache and kCausalConv1dFwd, both host-side by design. WHY THIS OP. After the GDN rows there was no kernel-speed lever left: both decode GEMMs sit near the GB10 bandwidth roof. What remained was the reference tier's ARCHITECTURAL cost -- src/vt/op_provider.cpp drains the recorded command batch (submit plus blocking fence) before it can hand a host kernel device memory, so a reference-tier op costs a full GPU round trip however little arithmetic it does. Of the three ops still declining, kCausalConv1dFwd is prefill-only and kRopeCosSinCache is deliberately on the host (its angle table is built in double precision, vLLM's own split), which left this one. The 27B has 64 layers of which 48 are linear-attention, so it fires 16x per token. PORTED FROM. Per-element math 1:1 from src/vt/cpu/cpu_ops.cpp:956-1010, our transcription of vLLM's QKNormRoPEFusionPass -> _C.fused_qk_norm_rope. Dispatch shape from src/vt/cuda/cuda_ops.cu:1316-1394: one workgroup per (token, head) over Hq + Hkv slots, CUDA's dim3(t, hq+hkv) flattened, the two paired normed elements recomputed per output rather than staged in shared memory -- CUDA's choice, kept so both devices compute the same expression. It is a fusion of two kernels that already work: the reduction is vt_rms_norm_gated.comp's, the cos/sin indexing is vt_rope_from_cache.comp's minus the positions indirection. MEASURED, GB10, 1 prompt, 32-in, c1, page cache dropped before every run. Two-length flush diff (output-len 4 vs 12), both arms measured: per decoded token main this row reference-tier drains 16 0 command-buffer flushes 18 4 GPU dispatches 884 900 GPU-active 236.3 ms 234.6 ms wall (median TPOT) 253.5 ms 246.2 ms The GPU does the same work to within noise and takes four MORE dispatches; what changes is that 16 submit-plus-blocking-fence round trips per token stop happening. Decode over 6 order-alternated AB/BA pairs: this row wins 5 of 6, median TPOT 249.74 -> 242.34 ms = 4.00 -> 4.13 tok/s. Discarding the two pairs carrying an outlier leg leaves 4 of 4 and the same ~3%. llama.cpp Vulkan on this box is 4.35. Attributing the ~3% to the removed round trips is INFERRED: host time is wall minus GPU-active, a derived quantity, not directly instrumented. Gates re-run independently rather than taken on report: on llvmpipe from a clean build, test_vulkan_backend 29/29 (1786 assertions), test_opt_paged_engine on Vulkan 6/6 token-exact (96/96 tokens) with 0 declines, test_backend_cross_device 11/11, and gen-vulkan-spirv.py --check reproduces the committed SPIR-V byte-for-byte under the pinned glslang 16.5.0. CI's own runners pass build-test-vulkan, build-test-cpu and both sanitizers. The f32 arm of the new kernel is NMSE-tier, not bit-exact -- the workgroup tree reduction reorders the mean square, the same trade vt_rms_norm and vt_rms_norm_gated already make; the bf16-q/k arm the model actually uses IS bit-exact. ALSO IN THIS MERGE, and deliberately not silent: docs/BENCHMARKS.md is a keyed table, and #154 had appended a dated H2 section to it. That section plus the em-dash it carried was failing check-public-doc-tables on MAIN, and therefore the agent-record job on every open PR rather than only on the change that caused it. roll-benchmark-record.py --apply moves it verbatim into the append-only record. TWO GATES REMAIN RED AND ARE NAMED RATHER THAN PAPERED OVER. check-pr-size reports the product class at 1306 lines against a 900 budget; 723 of those are the regenerated src/vt/vulkan/vulkan_spirv.cpp, a machine-generated hex blob whose freshness the --check gate proves, so hand-written lines are ~583, inside budget. The checker classifying generated SPIR-V as product is a real gap, left open. check-public-doc-tables still fails on docs/STATUS.md at 277213 chars against a 276960 shrink-only ratchet: this branch SHRINKS that page by 4 chars and main is already 257 over, and the only block big enough to clear it is a single 33,211-char table cell that is itself far past the 220-char cell rule. Collapsing a cell that size is a separately-reviewable change with its own owner. README.md:310 still reads "24 native ops"; check-doc-checkpoint rejects a README change without an accompanying landing-page source, so it needs to ride a change that qualifies. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…lags main has been red on every run overnight. Four of the failures are TREE-scoped, which is the part that matters: unlike the diff-scoped range gates, a tree-scoped failure reds every run AND every open PR until it is fixed, so it blocked unrelated contributor work. 1) check-public-doc-tables: docs/STATUS.md was 277214 chars against a 276960 shrink-only ratchet. The DeepSeek last-mile Bricks 3 and 4 were a run-by-run log of a NOT-pushed branch, 3871 chars for two negative results, on a page whose contract is one binding line per capability. Collapsed to 1009 chars that KEEP both negatives and the finding they establish: dp4a near-flat +0.5%, aligned repack a MEASURED NEGATIVE at +4.7 GiB, coalesced-load REFUTED, the Q8_0 matvec LATENCY/OCCUPANCY- bound at ~61% of ds4. The nsys tables were already in .agents/benchmark-record.md. Ratchet lowered 276960 -> 274760. The 2026-08-08 ratchet left ~15 chars of headroom. That was mine and it was too tight: three Vulkan landings that night each added a status line and pushed the page over within hours. The new headroom is 408, so the ratchet measures bloat rather than merge timing. 2) check-env-doc: 12 production env vars were undocumented -- seven VT_GEMMA4_*, three VT_ROCM_*, and the two VT_SERVER_MAX_* request limits. Documented in docs/ENVIRONMENT.md next to their siblings rather than allowlisted, because the existing VT_GEMMA4_RESIDENT_* and VT_ROCM_ATTN_CPU_REF knobs are documented there. Defaults were read from the source, not guessed. 3) check-state-order: the 2026-08-09T13:00 anchor sat before its heading instead of on the line after it, so that entry counted as unanchored. 4) check-now-current: NOW.md was 6080 chars against a 6000 budget. Dropped the two rows carrying no open work -- MXFP4 parity, marked TERMINAL, and the supported-models list, marked LANDED with a "-" residual. Both keep their evidence in docs/FEATURES.md and the matrices. Also the docs #154 owed, which is what reddened documentation-checkpoint when it landed. It shipped eight real user-facing flags with no USAGE entry: --repeat on vllm-cli, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose on the server. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. README: NOT changed here, and that is the one thing this PR could not fix. The count says 24 native ops in two places where STATUS and NOW both say 25, which .agents/state.md had already flagged. POL-DOC-README only accepts a README edit alongside one of six landing-source files (.agents/mission.md, CMakeLists.txt, benchmarks/demo/*.json, examples/{cli,server}/main.cpp), and a stale-number correction touches none of them and should not have to. Left for the developer to rule on, as state.md asked: either waive it or let the gate accept a correction whose source already sits in docs/BENCHMARKS.md. 5) tests/scripts/test_check_public_doc_tables.py had its `if __name__ == "__main__": unittest.main()` block at line 362, BEFORE the StatusRatchet class at line 392. CI runs this file as a script, so unittest.main() executed before that class was defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never ran. The ratchet's own mutation suite was inert. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass. Found because check-pr-size requires mutation evidence for a checker change and the new test I added did not fail when it should have. It was not failing because it was not running. The new test pins what this change actually alters: ratchet headroom is bounded, so a red char ratchet cannot be cleared by inflating the number instead of shrinking the page. Two-way mutation proof, bytecode disabled: chars=279000 -> FAILED "4648 not less than or equal to 2000"; chars=274000 -> FAILED "-352 not greater than or equal to 0"; baseline 274760 -> OK. Not fixed here, because it is not a bug: POL-PR-REQUIRED keeps firing on content commits from --merge landings and direct pushes to main. That gate is working. The permanent fix is the repository setting (allow squash only), which needs admin. Gates: all 20 tree-scoped checkers green, plus test_agent_record, test_doc_checkpoint, test_check_public_doc_tables, test_check_readme_structure, test_check_env_doc, test_check_state_order, test_check_now_current, test_agent_role, test_agent_gates, test_check_gate_commands and test_audit_live_rows -- on a worktree pinned at f921062. No code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…lags main has been red on every run overnight. Four of the failures are TREE-scoped, which is the part that matters: unlike the diff-scoped range gates, a tree-scoped failure reds every run AND every open PR until it is fixed, so it blocked unrelated contributor work. 1) check-public-doc-tables: docs/STATUS.md was 277214 chars against a 276960 shrink-only ratchet. The DeepSeek last-mile Bricks 3 and 4 were a run-by-run log of a NOT-pushed branch, 3871 chars for two negative results, on a page whose contract is one binding line per capability. Collapsed to 1009 chars that KEEP both negatives and the finding they establish: dp4a near-flat +0.5%, aligned repack a MEASURED NEGATIVE at +4.7 GiB, coalesced-load REFUTED, the Q8_0 matvec LATENCY/OCCUPANCY- bound at ~61% of ds4. The nsys tables were already in .agents/benchmark-record.md. Ratchet lowered 276960 -> 274760. The 2026-08-08 ratchet left ~15 chars of headroom. That was mine and it was too tight: three Vulkan landings that night each added a status line and pushed the page over within hours. The new headroom is 408, so the ratchet measures bloat rather than merge timing. 2) check-env-doc: 12 production env vars were undocumented -- seven VT_GEMMA4_*, three VT_ROCM_*, and the two VT_SERVER_MAX_* request limits. Documented in docs/ENVIRONMENT.md next to their siblings rather than allowlisted, because the existing VT_GEMMA4_RESIDENT_* and VT_ROCM_ATTN_CPU_REF knobs are documented there. Defaults were read from the source, not guessed. 3) check-state-order: the 2026-08-09T13:00 anchor sat before its heading instead of on the line after it, so that entry counted as unanchored. 4) check-now-current: NOW.md was 6080 chars against a 6000 budget. Dropped the two rows carrying no open work -- MXFP4 parity, marked TERMINAL, and the supported-models list, marked LANDED with a "-" residual. Both keep their evidence in docs/FEATURES.md and the matrices. Also the docs #154 owed, which is what reddened documentation-checkpoint when it landed. It shipped eight real user-facing flags with no USAGE entry: --repeat on vllm-cli, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose on the server. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. README: NOT changed here, and that is the one thing this PR could not fix. The count says 24 native ops in two places where STATUS and NOW both say 25, which .agents/state.md had already flagged. POL-DOC-README only accepts a README edit alongside one of six landing-source files (.agents/mission.md, CMakeLists.txt, benchmarks/demo/*.json, examples/{cli,server}/main.cpp), and a stale-number correction touches none of them and should not have to. Left for the developer to rule on, as state.md asked: either waive it or let the gate accept a correction whose source already sits in docs/BENCHMARKS.md. 5) tests/scripts/test_check_public_doc_tables.py had its `if __name__ == "__main__": unittest.main()` block at line 362, BEFORE the StatusRatchet class at line 392. CI runs this file as a script, so unittest.main() executed before that class was defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never ran. The ratchet's own mutation suite was inert. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass. Found because check-pr-size requires mutation evidence for a checker change and the new test I added did not fail when it should have. It was not failing because it was not running. The new test pins what this change actually alters: ratchet headroom is bounded, so a red char ratchet cannot be cleared by inflating the number instead of shrinking the page. Two-way mutation proof, bytecode disabled: chars=279000 -> FAILED "4648 not less than or equal to 2000"; chars=274000 -> FAILED "-352 not greater than or equal to 0"; baseline 274760 -> OK. 6) check-pr-size classified neither CLAUDE.md (added by the b8d996e direct push) nor MANIFESTO.md (a9a8581), so classify_path RAISED on them and any PR touching either failed the gate outright. Both are top-level reader-facing prose: same class as README. Mutation: dropping CLAUDE.md back out gives "Lists differ: ['CLAUDE.md'] != []". Not fixed here, because it is not a bug: POL-PR-REQUIRED keeps firing on content commits from --merge landings and direct pushes to main. That gate is working. The permanent fix is the repository setting (allow squash only), which needs admin. Gates: all 20 tree-scoped checkers green, plus test_agent_record, test_doc_checkpoint, test_check_public_doc_tables, test_check_readme_structure, test_check_env_doc, test_check_state_order, test_check_now_current, test_agent_role, test_agent_gates, test_check_gate_commands and test_audit_live_rows -- on a worktree pinned at f921062. No code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…#154 docs #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…#154 docs #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…#154 docs (#188) #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code] Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
…#154 docs #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. 7. check-device-leakage went red on main too, a fifth tree-scoped gate: DSR bucket vt_ifdef 37 > baseline 32. Cause is mechanical -- #189 moved the server body from examples/server/main.cpp, which the scanner never looked at, into src/vllm/entrypoints/openai/server_main.cpp, which IS the shared layer, carrying its 5 `#ifdef VT_BENCH_PROFILE_CONTROL` sites with it. The code did not change. Exempted per-site with the checker's own `// DSR-ALLOW(<row-id>):` hatch rather than a per-file ALLOWLIST budget, for two reasons. A budget can be spent on something else -- swap one guard for a real device fork and the count still reads 5 -- while a per-site marker names each one. And the ALLOWLIST route required editing check-device-leakage.py, which check-pr-size then rejected with "has no affected POL rule mapping": that checker is named by no POL rule in policy.csv, so registering it is a policy decision, not a CI repair. The baseline is NOT raised; DSR returns to 32 == 32 and the 5 exemptions print in CI output every run. The real repair is to put the profiler seam behind a Platform capability query like every other device fork, which is a change to #189's TU. 8. check-pr-size could never accept a change to check-device-leakage.py. recognized_evidence derives tests/scripts/test_check_<name>.py, but that suite predates the convention and CI runs it as tests/scripts/test_device_leakage.py, so every change to that checker failed with "requires semantic mutation evidence in <a file that does not exist>" -- which is exactly what happened above. Added to CHECKER_EVIDENCE_OVERRIDES, which exists for this. An audit of all 30 checkers found one more: check-dsv4-gguf-namemap.py has no suite at all and is not in ci.yml, named in a KNOWN_UNTESTED set so the gap is visible in a test rather than invisible in a naming rule. Mutations: removing the override gives "Lists differ: ['scripts/check-device-leakage.py -> ...'] != []"; downgrading one DSR-ALLOW to a plain comment fails both the new per-site test ("line 50 has no DSR-ALLOW on it or directly above it") and the DSR ratchet (33 > 32). Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
…#154 docs #190 landed while this branch was open and fixed more of main's red than this did, and did it better: the STATUS.md ratchet cleared by MOVING a 33,211-char cell verbatim per POL-EVIDENCE-PRESERVE rather than condensing it, CLAUDE.md/MANIFESTO.md classified (CLAUDE.md as procedure -- it is a symlink to AGENTS.md -- which is more accurate than the public_document I had used), a generated path class for SPIR-V, and a real benchmarks/demo/vulkan_27b_llamacpp.json landing source so the false README could finally be corrected. Everything this branch had that #190 also covers is dropped rather than re-litigated. What was still red on main at c1716fd, both TREE-scoped, so both were failing every run and every open PR: 1. check-agent-record: .agents/parity-ledger.md carried `../examples/server/main.cpp#L90`, and #189 turned that file into a 23-line thin ABI client, so the anchor pointed past EOF. The referenced "server flag" is the CUDA-graph-replay profiler trigger, which now lives at src/vllm/entrypoints/openai/server_main.cpp:692 (`args.cuda_profile_graph_replays > 0`). Repointed there. Only the pointer moves; the ledger entry is untouched. 2. check-env-doc: VLLM_CPP_HTTP_FIXED_POOL, added by #189, was undocumented. Documented in docs/ENVIRONMENT.md rather than allowlisted, because it changes runtime behavior (`=0` reverts the HTTP worker pool to legacy dynamic) and the comparable A/B switches VT_ROCM_ATTN_CPU_REF and VT_DEBUG_SAMPLED are documented there. 3. tests/scripts/test_check_public_doc_tables.py still has its `if __name__ == "__main__": unittest.main()` at line 362, BEFORE the StatusRatchet class at 392. CI runs this file as a script, so unittest.main() executes before that class is defined and all 8 of its tests -- including test_growth_past_the_char_ratchet_is_rejected and test_the_live_page_is_inside_its_ratchet -- never run. The ratchet's own mutation suite is inert, which is notable given #190 just re-pinned that ratchet. Moving the block to the end takes the file from 41 to 49 collected tests; all 49 pass against #190's 244486. 4. docs/USAGE.md: #154 shipped eight user-facing flags with no USAGE entry and #189 did not add them either. Re-verified every one still parses after the thin-client refactor: --repeat in examples/cli/main.cpp, and --enable/--disable-log-requests, --enable-log-outputs, --max-log-len, --enable/--disable-metrics, --enable/--no-enable-thinking and --verbose in src/vllm/entrypoints/openai/server_main.cpp. NOT documented: --default-chat-template-kwargs, which appears only in a source comment naming vLLM's spelling and is not a parsed flag. 5. docs/BENCHMARKS.md still told the reader the Vulkan-vs-llama.cpp arm was "Not yet runnable (no model runs on Vulkan)". It runs; that row now names the harness. 6. docs/STATUS.md said the multimodal server seam "is wired into `examples/server/main.cpp`", which #189 made false -- the `preprocessor_config.json` guard now lives in src/vllm/entrypoints/openai/server_main.cpp (verified by grep, not assumed). Corrected. The longer path costs 19 chars against 3 of ratchet headroom, so it is paid for inside the same cell by tightening two clauses that restate what the sentence already says. Net -5. 7. check-device-leakage went red on main too, a fifth tree-scoped gate: DSR bucket vt_ifdef 37 > baseline 32. Cause is mechanical -- #189 moved the server body from examples/server/main.cpp, which the scanner never looked at, into src/vllm/entrypoints/openai/server_main.cpp, which IS the shared layer, carrying its 5 `#ifdef VT_BENCH_PROFILE_CONTROL` sites with it. The code did not change. Exempted per-site with the checker's own `// DSR-ALLOW(<row-id>):` hatch rather than a per-file ALLOWLIST budget, for two reasons. A budget can be spent on something else -- swap one guard for a real device fork and the count still reads 5 -- while a per-site marker names each one. And the ALLOWLIST route required editing check-device-leakage.py, which check-pr-size then rejected with "has no affected POL rule mapping": that checker is named by no POL rule in policy.csv, so registering it is a policy decision, not a CI repair. The baseline is NOT raised; DSR returns to 32 == 32 and the 5 exemptions print in CI output every run. The real repair is to put the profiler seam behind a Platform capability query like every other device fork, which is a change to #189's TU. 8. check-pr-size could never accept a change to check-device-leakage.py. recognized_evidence derives tests/scripts/test_check_<name>.py, but that suite predates the convention and CI runs it as tests/scripts/test_device_leakage.py, so every change to that checker failed with "requires semantic mutation evidence in <a file that does not exist>" -- which is exactly what happened above. Added to CHECKER_EVIDENCE_OVERRIDES, which exists for this. An audit of all 30 checkers found one more: check-dsv4-gguf-namemap.py has no suite at all and is not in ci.yml, named in a KNOWN_UNTESTED set so the gap is visible in a test rather than invisible in a naming rule. Mutations: removing the override gives "Lists differ: ['scripts/check-device-leakage.py -> ...'] != []"; downgrading one DSR-ALLOW to a plain comment fails both the new per-site test ("line 50 has no DSR-ALLOW on it or directly above it") and the DSR ratchet (33 > 32). Note for whoever touches docs/STATUS.md next: #190 tightened the char ratchet to the measured byte, so slack is 3. Any status line added without paying for it reds a tree-scoped gate again. That is a deliberate choice of #190's, left alone here. Gates: 14 tree-scoped checkers green, plus test_check_public_doc_tables (49), test_agent_record, test_doc_checkpoint, test_check_state_order, test_check_now_current, test_check_env_doc and test_check_pr_size (21), on a worktree pinned at c1716fd. No product code touched. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude-Code:claude-opus-5 [Claude Code]
Summary
Follow-up to merged #140 (ROCm gfx1201 + Gemma-4-26B MoE BF16/FP8).
This PR hardens dual-GPU FP8 expert residency and decode mix so the path fits ~30 GB host RAM and keeps weights on GPU:
MulScalar+Addfor expert weighted sum (default;VT_GEMMA4_HOST_AXPY=1fallback)Lab evidence (2× R9700 gfx1201, ROCm 7.2.x)
<bos>The capital of France is→**Paris**EXIT=0Host OOM root cause was resident upload calling cache-filling dequant for every expert every layer.
Usage
HIP_VISIBLE_DEVICES=0,1 \ VT_GEMMA4_RESIDENT_EXPERTS=1 \ VT_GEMMA4_RESIDENT_GPUS=2 \ ./build-hip/examples/vllm-cli \ --model /path/to/Firworks/gemma-4-26B-A4B-it-fp8 \ --prompt '<bos>The capital of France is' --max-tokens 12 --temperature 0Optional:
VT_GEMMA4_HOST_AXPY=1,VT_GEMMA4_RESIDENT_MAX_LAYERS=NTest plan
Notes