diff --git a/docs/INDEX.md b/docs/INDEX.md index 65f2226ee..e03542bc3 100644 --- a/docs/INDEX.md +++ b/docs/INDEX.md @@ -105,6 +105,7 @@ Exactly one canonical owner (or explicit `BLOCKED`) per concern. | Relicense / governance records | [`docs/governance/`](governance/) | historical | Legal/historical; do not rewrite. | | Design-time architecture drafts | [`docs/design/`](design/) | planned | Directory policy is planned intent. Member exceptions: some branch-implemented LFM/Redline designs and measured baseline notes may coexist; none are product defaults. Historical member bodies stay disposition-only. | | Implementation plans and PRDs | [`docs/plans/`](plans/) | planned | Directory policy is planned intent. Member exceptions: measured results ledgers may appear; recency ≠ authority. | +| Codebase audit 2026-09-02 | [`docs/plans/audit-2026-09-02/`](plans/audit-2026-09-02/README.md) | planned | Point-in-time audit of master `8cd15a62b` (twelve read-only slices) with a ranked remediation plan. Findings cite `path:line` at that commit. | | Narrow specs | [`docs/specs/`](specs/) | planned | Intent/spec records; promote only via shipped owners. | | Investigations | [`docs/investigations/`](investigations/) | historical | Discovery trails including measured research; not product defaults. | | Reviews | [`docs/reviews/`](reviews/) | historical | Review archives. | diff --git a/docs/plans/audit-2026-09-02/README.md b/docs/plans/audit-2026-09-02/README.md new file mode 100644 index 000000000..e5d29e0bd --- /dev/null +++ b/docs/plans/audit-2026-09-02/README.md @@ -0,0 +1,72 @@ + + +# hipfire codebase audit — 2026-09-02 + +_Lifecycle: planned intent (docs/INDEX.md § plans). A point-in-time audit of master `8cd15a62b` with a ranked remediation plan; findings are cited to `path:line` at that commit and go stale as those lines move._ + +Master `8cd15a62b`. Twelve read-only scouts, one slice each, same rubric (broken / missing / would change, every finding cited to `path:line`). Their full reports are alongside this file as `audit-.md`. A follow-up kernel-family audit of MQ4G256V2 (qt=44), the production 4-bit format, is in [`audit-Mq4v2Kernels.md`](audit-Mq4v2Kernels.md) — verdict: sound end to end, one runtime admit out of lockstep, one non-discriminating parity fixture. A DFlash audit is in [`audit-Dflash.md`](audit-Dflash.md) — headline: DFlash is unreachable without an explicit draft path (the documented pull-and-`auto` on-ramp does not exist), draft VRAM (~2.5 GiB on 27B at the default cap) is charged nowhere, and one post-prefill error exit skips rollback. I spot-checked the five findings below marked ★ by reading the cited lines myself; the rest are the scouts' verified claims and should be read as "cited, plausible, not re-derived by me." + +## The shape of it + +The codebase is in better shape than the clobber week suggested, and the failure modes are consistent. Almost nothing here is a wrong computation. Almost everything is one of three things: + +1. **A door the loader opens that the engine can't walk through.** The loader admits a (model, topology, quant) combination, spends the VRAM, and then the generate side has no arm for it — or has an arm that refuses. #683 (Qwen3.5-MoE EP) was the first one found; the audit found five more of the same species. +2. **Two halves of a migration both alive.** A "canonical" path was introduced (`reset_lifecycle`, `production_fail_closed_rollback`, `ProcessConfig`, `ArchModel`) and most routes moved to it, but a handful of routes — usually the bring-up archs and the vision loops — still carry the old hand-rolled version, which has since drifted. +3. **Comments and docs that describe the design, not the code.** `SpecLoadCfg` says it falls back to ambient env; it doesn't. `gemm_table` says V2 is gfx12-only; it isn't. `env-vars.md` says production never reads ambient `HIPFIRE_*`; the daemon reads six. These aren't cosmetic — two of them (the vmm-on-EP comment in #682, the "mirrors" claim) are exactly how the clobber got approved. + +Nothing in the audit is a security problem. Nothing is data loss. The worst class is "a user configures a supported-looking thing, waits for a 19 GB load, and gets an error" — and the second-worst is "a reload leaks VRAM until the process dies." + +## Broken — the ones that matter, ranked + +**1. ★ Loader admits, engine refuses (the #683 family).** Six instances, all verified: +- Qwen3.5-MoE EP: `hipfire-generate/src/qwen.rs:240` routes arch 6 to the dense server (#683, already filed). +- **Gemma4 lowered/MoE**: the carrier loads `Gemma4Lowered` (`hipfire-arch-gemma4/src/carrier.rs`), and `ar.rs:1160-1170` hard-refuses generate on it — "not yet wired on this build". The load succeeds first. +- **Any EP arch other than 9/10 that reaches `generate_ep`** falls through to `ep_serve_ds4` (`ar.rs` `generate_ep` `_ =>` arm) — wrong server, not a refusal. LFM2/Cohere2 have no `EpArch` at all, so a forced EP load lands here. +- **LFM2 continuous batching**: `batch.rs:166-181` runs four eligibility checks and then `return false` unconditionally. `caps.supports_continuous_batch` is `true` for LFM; the batch state gets allocated and never used. +- **Muse-Glimmer continuous batching**: same pattern — `GlimmerDecodeBatchState` exists (`hipfire-arch-muse-glimmer/src/batch.rs`), `continuous_batch_route` only admits `5|6|11`. +- **GL quant types 38/39** load via `RAW_CODECS` (`weight_backend.rs`) and have no dense GEMV dispatch (`hipfire-dispatch/src/types.rs`), so a GL-quantized dense model loads and cannot decode. + +Fix shape is the same for all six: refuse at admission with the reason, before allocation. That's a one-arm change each and the cheapest correctness win in the codebase. Wiring the features is separate work. + +**2. ★ MQ8 GEMV family path panics.** `hipfire-dispatch/src/families/gemv.rs:140-148` maps `RotationPlan::Mq8Internal` to a plain FWHT rotation, which never fills `scratch.mq_x_q8`; `rdna-compute/src/gemv.rs:6163` then `unwrap()`s it. The pipeline path (`steps.rs:940`) special-cases MQ8 correctly and is what production uses today, so this is a landmine on the `run_auto` path rather than a live crash — but the plain `KernelKey::GemvMq8G256` is registered and falls through to `MissingImpl` too, so MQ8 has exactly one working route and two broken ones. Hours to fix; add a unit test that runs every registered gemv key once. + +**3. ★ VRAM leaks on reload.** `LlamaWeights::free_gpu` (`hipfire-runtime/src/llama.rs:683`) frees `.buf` per weight and skips the PARO rotation and AWQ scale sidecars; the leak repeats every reload of an AWQ/PARO llama/qwen3. `DeviceBuffer` has no `Drop` (`hip-bridge/src/lib.rs`), so every owner that forgets `free` leaks silently — the audit found the `dspark` weights, the DFlash layer staging on partial failure, and the `Qwen3DsparkScratch` intermediates all in that state. (Several of these were what the reverted G4 series *did* fix; those small pieces are worth re-landing individually.) + +**4. Reset/rollback is canonical on the main routes and hand-rolled everywhere else.** `production_fail_closed_rollback` covers AR, dense-TP, and DS4. Verified gaps: EP GPU-error exits skip `ep_reset_after_abort` (multi-rank state left dirty — MiniMax LCP especially); the vision cancel path emits a wire terminal without a GPU/DeltaNet rollback (`vision.rs`); vision context-full still does a manual DN memset; the bring-up dense archs (qwen2, minimax, maple, cohere) have an `Abort => {}` arm — no rollback at all. `reset_core_arch_key` maps arch 15 (maple) to `"unknown"`, and `dots_ocr` vs `dots-ocr` hyphen drift means dots never gets reset-core coverage. Any of these is a next-turn corruption bug on a cancel or an error. + +**5. EP loads build a Tp mesh.** ★ Every `load_model_ep_*` calls `Gpus::init_tp` (`hipfire-loader/src/lib.rs:2994,3222,3361`), which records `DeviceMesh::rect(Tp, n)` (`multi_gpu.rs:300`). So after an EP load, `mesh.size_of(Ep) == 1`. `#681` just landed this type with no readers, so it's harmless today — but the first consumer of the mesh will get the wrong topology for every MoE load. `Gpus::init_ep` is an hour of work and should land before anyone reads `mesh`. Also: `Gpus::single` records an empty mesh while `from_parts` records `Pp:N` for N=1 — two representations of the same thing. + +**6. Config bypasses.** `docs/env-vars.md:51-69` says production never reads ambient `HIPFIRE_*` after `ProcessConfig` is installed. The daemon reads `LOG_FORMAT`, `DFLASH_DRAFT`, `PP_DFLASH`, `PFLASH`, `DPM_WARMUP` directly (`main.rs:497,1199,1575-1589,1801`), the loader reads `PAGE_EVICTION` (`carriers.rs:409`), and `hipfire bench` sets `continuous_batch` via `set_var` as a side channel to itself (`cli/main.rs:3925`). `check-env-docs.py` only catches literal `HIPFIRE_*` string reads and reports success on this tree. The PP experimental gates are snapshotted into config at startup *and* re-read live — two sources of truth for the same flag. + +**7. Docs contradict the build.** `CONTRIBUTING.md:58`, `GETTING_STARTED.md:129`, `VALIDATION.md:84-97` still say `cargo run --example daemon -p hipfire-runtime` and reference a `test_kernels` example that doesn't exist; the daemon has been a crate since the saddle. `CONTRIBUTING.md:33-44` links `.skills/`; the real path is `.agents/skills/`. `ARCHITECTURE.md:107` says `--json`/`--no-stream` force local; `main.rs:1911-1918` shows they don't. `MODELS.md` says deepseek-v4-flash defaults to `mq2lloyd`; the registry says `mq2r`. README's model count is three different numbers. `CONTRIBUTING.md:286-295` asks contributors to help with issues #57 and #58, both closed in April/May. + +**8. Smaller verified items** (each an hour): `gemm_table.rs:364` comment says V2 is gfx12-only while the code gates on `HasWmma` and gfx11 sources exist; MQ6 residual `ensure_kernel` uses the `_mq5v2` module suffix (`gemm.rs:32097`, no collision, wrong name); Cohere2-MoE docs say LayerNorm, forward does RMSNorm; `SpecLoadCfg` doc claims an env fallback that `spec_build.rs:201` doesn't implement; `verify-bind-thread.sh` only audits `dispatch.rs`, while `impl Gpu` is split across ten files; `hsa-bridge` has ~30 `unsafe` blocks with no SAFETY comments; the registry silently drops aliases whose target doesn't exist (`hipfire-registry/src/lib.rs:328`). + +## Missing + +- **`Gpus::init_ep`** and EP-mesh admission (see 5). +- **A `DeviceBuffer` Drop or a `#[must_use]` free** (see 3) — the one structural change that would retire a whole bug class; days, because VMM-backed buffers have different owners. +- **Fail-closed `generate_ep`** and EP admission refusal for arch 11/12/13. +- **Docs-drift CI** beyond `check-env-docs.py` — the entrypoint/path/count drift above would all be caught by a script that greps docs for `cargo run --example`, `.skills/`, and compares README counts to `registry/models.json`. +- **Windows device selection** (issue #669): `hardware.devices` → `ROCR_VISIBLE_DEVICES` doesn't work there; no docs, no fallback. HUSRCF's workaround is `CUDA_VISIBLE_DEVICES`, which nothing documents. +- **LFM2 multi-turn**: always cold-resets conv+KV (`dense.rs:6105`), so every turn re-prefills. Correct, but it's why LFM feels slow in chat; a conv-state snapshot is a few days. +- **DS4 `route_scale`** is forced to 2.2/1.8 instead of the checkpoint's 1.5 (`config_cache.rs:450-510`) to compensate for a MoE routed-branch shortfall nobody has root-caused. It works; it's also a hidden numeric fudge that will bite the next quant. +- `hipfire registry` subcommand is absent from `docs/CLI.md`; `init_vram_weighted` is a stub that errors; `ds4-parent` is an offline oracle and should stay quarantined (it is). + +## Would change — what I'd actually do, in order + +1. **Admission refusals for the six #683-class combinations** (hours, one PR, ships behind hw-gate). Biggest user-facing win per hour in the codebase; every one currently costs a user a full load before the error. +2. **`Gpus::init_ep` + unify `single`'s mesh** (hours). Land before anything reads `mesh`. +3. **Re-land the G4 leak fixes as individual PRs**: `LlamaWeights` sidecar frees, DFlash/DSpark staging on failure, `Qwen3DsparkScratch` (days total, each one a small PR with a fault scenario, gated). +4. **Close the reset gaps**: EP error exits → `ep_reset_after_abort`; vision cancel → rollback; bring-up `Abort => {}` arms → the canonical rollback; fix `reset_core_arch_key` for maple and the dots hyphen (days). This is the G4 work that was actually right, done at a size that can be reviewed. +5. **MQ8 rotation family** (hours) + a test that launches every registered gemv key once. +6. **Route the residual ambient env reads through `ProcessConfig`** and make `check-env-docs.py` fail on them (1–2 days). Until then `env-vars.md`'s promise is false. +7. **Docs pass** (a day): entrypoints, `.agents/skills`, force-local, model counts, MODELS default, ARCHITECTURE carrier table, INDEX pin, closed-issue asks. Then a drift script in CI. +8. **`DeviceBuffer` RAII** (days, coordinate VMM). The structural fix for class 3. +9. **Issue triage** (a day): 88 open; the scout's table in `audit-DocsIssues.md` clusters them. #675/#677/#527 are historical after the revert; #593/#595/#563 are on a stale beta base; ~20 have had no response in 30+ days. + +Things I would explicitly *not* change now: the DS4 route-scale compensation (works; root-cause first), PFlash (retired, quarantined, not reachable from the daemon), `ds4-parent`. + +## Confidence + +Twelve slices, ~7 minutes each, read-only. The scouts were told to prefer 10–20 findings that matter over exhaustiveness, so the long tail of the arch crates (deepseek4's 36k lines, rdna-compute's 178k) got a survey, not a read. Kernels (`kernels/`, 146k lines HIP) were sampled for dispatch coverage only; nothing in this audit says the math is right — that's what the redline/parity routes are for. I re-derived five findings by reading the lines; the rest carry the scouts' citations. No finding contradicts another scout's, and the three that overlap (#683 family, reset gaps, VRAM leaks) were found independently by two or three slices each, which is the best signal that they're real. diff --git a/docs/plans/audit-2026-09-02/audit-ArchMoE.md b/docs/plans/audit-2026-09-02/audit-ArchMoE.md new file mode 100644 index 000000000..21185db52 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-ArchMoE.md @@ -0,0 +1,96 @@ + + +# Audit: ArchMoE + +# ArchMoE audit — hipfire origin/master @ 8cd15a62b + +Scope: hipfire-arch-deepseek4, minimax, cohere2moe, lfm2moe, lfm2-vl, ds4-parent + generate/loader EP wiring. Read-only. gh issue keyword scan MoE/EP/DSML: no matching open issues (exit 0). Arch IDs: DS4=9, MiniMax=10, LFM2-MoE=11, Cohere2-MoE=12; LFM2-VL vision for 11. + +## Broken + +### 1. EP Unknown fallthrough runs DeepSeek4 EP server (high) +**path:** crates/hipfire-generate/src/ar.rs — select_generation_route + generate_ep +**verified:** yes + +EP arms only 9→Deepseek4Ep, 10→MiniMaxEp; else Unknown (incl. Qwen 5|6 EP, future 11/12). generate_ep: 10→ep_serve_minimax, 5|6→ep_serve_qwen35_dense_tp, **_→ep_serve_ds4**. Any Unknown EP reaching generate_ep runs DS4 EP against wrong weights/protocol (#683 class). Loader today only EpArch::Ds4|Minimax limits blast radius; match still fail-open. + +### 2. LFM continuous-batch admission always false (high for dense LFM) +**path:** crates/hipfire-generate/src/batch.rs:168-180 +**verified:** yes + +After LfmAr, lfm2_decode_batch, is_dense(), batch_weight_formats_supported → unconditional `return false`. carriers.rs:1470 supports_continuous_batch:true; staging can build dense batch — unreachable. MoE LFM intentionally excluded (batch.rs~84-86); dense LFM was meant to pass. Dead path / half-migration. + +### 3. Cohere2-MoE docs vs norm implementation +**path:** hipfire-arch-cohere2moe lib.rs/map vs forward RMSNorm +**verified:** yes + +Docs: mean-centered LayerNorm. Code: rmsnorm_batched / RMSNorm at rms_norm_eps. Code/docs disagree (stale docs vs wrong math). + +### 4. DS4 routed MoE scale is permanent defect compensation +**path:** crates/hipfire-arch-deepseek4/src/config_cache.rs:450-510 resolve_route_scale +**verified:** yes + +Ignores cfg.routed_scaling_factor (checkpoint ~1.5). Defaults 2.2 non-mq2r / 1.8 mq2r; env HIPFIRE_DEEPSEEK4_ROUTE_SCALE wins. Restoring 1.5 cost ~51% PPL (16.31 vs 10.81 ctx2048); reference fine at 1.5 (PPL~4.7) — systematic hipfire MoE routed-branch shortfall. Parent path must keep 1.5. + +### 5. MTP/SWA reject stale ring residual +**path:** crates/hipfire-arch-deepseek4/src/spec_decode.rs (+ spec_impl/MTP) +**verified:** yes + +Rejected drafts can leave position-indexed SWA/DSA rings inconsistent unless rollback mirrors dense partial-LCP cold-rebuild. Cross-turn reset pairs state.reset + zero_decode_caches; intra-step partial reject weaker. In-code production-hardening follow-up. + +## Missing + +### 1. LFM2-MoE and Cohere2-MoE EP serve +**path:** ar.rs EP arms; no EpArch 11|12 +**verified:** yes — only DS4/MiniMax EP; LFM supports_ep_batch false; forced EP → Unknown → DS4 fallthrough risk. + +### 2. Heterogeneous DS4: no spec/tools until G6 +**path:** carriers.rs:1127-1128; deepseek4 carrier ~76-77 +**verified:** yes — SpecTarget "direct-AR only until G6"; draft/DSpark refused until G6/G7. Dense DS4 keeps MTP/DSpark. + +Ownership (healthy): hetero Drop/release frees prefill→state→weights.free_gpu(dense,routed) then both pools; RoutedWeights::free_gpu only free_routed_gpu; MiniMax dummy_gate_up owned on layer (no mem::forget leak). + +### 3. LFM continuous batch product gap +MoE excluded; dense admission broken → no live LFM continuous batch despite caps. + +### 4. LFM multi-turn prefix reuse +**path:** dense.rs:6105-6121; lfm2moe state.reset memsets conv_states +**verified:** yes — every request cold-resets KV+conv; hybrid conv cannot cheaply rewind. Full history re-prefill each turn. LFM2-VL vision-only crate; text in lfm2moe. + +### 5. Cohere block-parallel speculative verify +**path:** spec_impl.rs:16-38 — sequential decode_step only; windowed causal-within-block FOLLOW-UP. + +### 6. Runtime expert paging +Load-time EP shard + hetero split only; no ExpertPager; experts process-lifetime resident. + +### DSML / invented tokens — clean +**path:** dsml.rs:39-57; git a4630b753 reverts 2b86c6f6e +a4630b753 restored HF DSML string markup after native-token experiment. String constants only; no hardcoded invented special-token ID table. Fail-closed Malformed on unclosed tool spans. Do not reintroduce numeric invented IDs. + +### ds4-parent reachability +**path:** hipfire-ds4-parent; grep loader/generate/daemon clean +Offline parent-checkpoint oracle/examples only. Banner: NOT production quant calibration ref (PPL~59 vs teacher~4.7). Must not drive serving route_scale or quant. Heterogeneous DS4 architecture-local; no parent hook in Qwen loaders. + +## Would change (ranked) + +1. Fail-closed generate_ep — explicit DS4 arm; Unknown→error; sync route table. **hours.** Highest priority. +2. Fix or delete LFM batch admission — remove return false OR caps false + delete staging. **hours.** +3. Loader refuse --ep for arch 11/12 (no EpArch). **hours.** +4. Cohere norm doc/code alignment (+ contract test). **hours.** +5. Root-cause DS4 routed MoE shortfall; keep compensation until fixed; parent stays 1.5. **week+.** +6. MTP reject → SWA/DSA scrub to accept boundary; parity vs AR. **days.** +7. LFM multi-turn cache only with conv snapshot proof — else document re-prefill permanent. **days** if pursued. +8. Quarantine ds4-parent in CI/docs (already out of serve). **hours.** + +## Confidence + +Did: crate maps; EP route/serve; DS4 ep+hetero free/Drop; MiniMax EP ownership; LFM reset/free + generate cold reset; LFM batch admission; Cohere spec_impl + RMSNorm vs docs; DSML strings + revert history; ds4-parent isolation; route_scale; carriers caps/G6; open issue scan empty. + +Did not fully deep-read: every MTP accept/reject block; full DSpark rollback matrix; all MiniMax/LFM spec_impl EP interaction; load_model_ep full body (inferred EpArch sites); hardware (forbidden). + +Suspicious not Broken: whether HTTP can set m.ep for non-9/10 today (likely no); whether Cohere RMSNorm matches shipped checkpoint. + +Peer boundaries: Generate/Runtime/ConfigTopology/Other own overlapping surfaces; ar.rs/batch.rs/carriers.rs cited for MoE wiring only. + +--- +CONTRACT_JSON_FOR_PARENT: use summary field embedded JSON (slice/broken/missing/changes); full markdown is this architecture field. report path would have been local://audit-ArchMoE.md — parent persists. diff --git a/docs/plans/audit-2026-09-02/audit-ArchOther.md b/docs/plans/audit-2026-09-02/audit-ArchOther.md new file mode 100644 index 000000000..eeeb06a38 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-ArchOther.md @@ -0,0 +1,84 @@ + + +# Audit: ArchOther + +# Audit ArchOther + +## Broken + +1. **Gemma4 lowered/MoE loads; generate hard-refuses (verified)** — `hipfire-generate/src/ar.rs:1160-1170`. If `gemma4_lowered_mut().is_some()`, emit *"lowered/MoE generate not yet wired (eager dense only)"*. Carrier dual Eager/Lowered (`hipfire-arch-gemma4/src/carrier.rs`). Load-OK/generate-fail archetype. Open #678 describes lowered generate on a branch; **master still refuses**. + +2. **Gemma4 EAGLE full code, not product path (verified)** — `hipfire-arch-gemma4/src/speculative.rs:1-50` `spec_step_gemma4_eagle`; drafter + `infer_gemma4_spec`; `dense.rs` ~2226 EAGLE arm when arch-22 drafter + greedy. Gate-off class a72279cb1 (greedy parity). Dual maintenance. + +3. **Optional bias panics (verified)** — `tools/change_gate/routes.py` documents `tensor not found: layers.0.self_attn.q_proj.bias`. `WeightBackend::bias` non-optional. 31B-class optional bias; 12B dense OK. Distinct from #614 hipGraph and #678 ThoughtRouter. + +4. **Muse-glimmer continuous batch implemented, never staged (verified)** — `muse-glimmer/src/batch.rs` `GlimmerDecodeBatchState` + `batch_weight_formats_supported`. `hipfire-loader/src/lib.rs:186-198` `continuous_batch_route` only 5|6|11. `carriers.rs:2123-2124` `supports_continuous_batch: false`. + +5. **MuseGlimmerBundle/ArchModel orphaned in loader (verified)** — no arch-crate `carrier.rs`; `MuseGlimmerCarrier`+bundle in `hipfire-loader/src/carriers.rs` arch 14. Contradicts `hipfire-runtime/src/arch.rs:1-30` (ArchModel in arch crate). + +6. **dots-ocr docs say load unwired; code wired (verified)** — `lib.rs`/`map.md` vs `carrier.rs`+Architecture+ArchModel+SpecTarget arch 8. + +7. **`Architecture::eos_filter_overrides` dead for product generate (verified)** — trait `arch.rs:191-194`; dots/gemma4 overrides; generate never calls them (Architecture used for cohere2moe/deepseek4/minimax load). `dense.rs:7678-7682` Maple no EosFilter. Dots tests only `arch.rs:138-142`. + +8. **DotsOcr SpecTarget vision lifetime contradiction (verified)** — `spec_impl.rs:58-68` comment vision NOT included/freed after prefill; struct holds full `DotsOcrWeights`. + +## Missing + +1. Gemma4 vision stubs (`gemma4_vision.rs`). +2. Maple: no Architecture/SpecTarget/batch; no EOS/think filter on generate. +3. Glimmer spec/MTP: drafter exists; caps MTP false; emitter unwired. +4. Toy unshippable (`0xFF`, load always Err) still sole “template” teaching Architecture-centric world. +5. saddle-core: caps real; Architecture naming split confuses; spec helpers thin. +6. No load-time refuse when generate cannot serve lowered gemma4. + +## Would change (ranked) + +1. Fail-closed gemma4 lowered at load **or** wire generate_lowered — carrier+ar.rs — **hours**. +2. Optional bias → Option/skip-missing — weight_backend+gemma4 — **hours–days**. +3. Fence/delete EAGLE product surface — speculative/drafter/dense — **days**. +4. Move MuseGlimmer Carrier+Bundle+ArchModel into arch crate — **days**. +5. Wire or delete glimmer continuous batch — batch+route+caps — **days–week+**. +6. dots-ocr doc cutover + vision lifetime + generate-side EOS — **hours**. +7. Maple EOS/stop on generate without full Architecture — **hours**. +8. Retire toy as sole template; point at maple+dots — **hours**. + +## Confidence + +Read-only master `/home/kaden/ClaudeCode/warpfront/hipfire`. No GPU/tests/builds. Did not exhaust every forward kernel body. saddle-quant → AuditQuantize; MoE/Qwen generate → peers. GitHub open: #614, #678, #672 not re-filed; findings additive. Re-check ar.rs if lowered generate merges. + +## Contract JSON (for Main persist) + +```json +{ + "slice": "ArchOther", + "broken": [ + {"title": "Gemma4 lowered/MoE loads but generate hard-refuses", "path_line": "hipfire-generate/src/ar.rs:1160-1170", "verified": true, "summary": "Loader publishes Gemma4Lowered; ar.rs refuses generate eager-only."}, + {"title": "Gemma4 EAGLE full code production-gated", "path_line": "hipfire-arch-gemma4/src/speculative.rs:1-50", "verified": true, "summary": "spec_step_gemma4_eagle+drafter live; product path gated."}, + {"title": "WeightBackend bias panics optional missing", "path_line": "tools/change_gate/routes.py", "verified": true, "summary": "q_proj.bias tensor not found panic; 31B-class."}, + {"title": "Glimmer CB implemented never staged", "path_line": "hipfire-loader/src/lib.rs:186-198", "verified": true, "summary": "batch.rs exists; route only 5|6|11; caps false."}, + {"title": "MuseGlimmer Bundle/ArchModel in loader", "path_line": "hipfire-loader/src/carriers.rs", "verified": true, "summary": "Orphan rule; no arch carrier.rs."}, + {"title": "dots-ocr docs claim load unwired", "path_line": "hipfire-arch-dots-ocr/src/lib.rs", "verified": true, "summary": "Carrier wired; docs stale."}, + {"title": "eos_filter_overrides not used by generate", "path_line": "hipfire-runtime/src/arch.rs:191-194", "verified": true, "summary": "Trait docs lie relative to generate."}, + {"title": "DotsOcr vision doc vs bundle", "path_line": "hipfire-arch-dots-ocr/src/spec_impl.rs:58-68", "verified": true, "summary": "Comment no vision; struct has DotsOcrWeights."} + ], + "missing": [ + {"title": "Gemma4 vision", "path_line": "hipfire-arch-gemma4/src/gemma4_vision.rs", "verified": true, "summary": "Stubs only."}, + {"title": "Maple Architecture/SpecTarget/EOS", "path_line": "hipfire-arch-maple/src/carrier.rs", "verified": true, "summary": "Thin surface; generate EOS gap."}, + {"title": "Glimmer spec/MTP", "path_line": "hipfire-loader/src/carriers.rs caps", "verified": true, "summary": "Drafter present; MTP false."}, + {"title": "Toy unshippable template", "path_line": "hipfire-arch-toy", "verified": true, "summary": "0xFF always-Err load."}, + {"title": "saddle-core Architecture naming split", "path_line": "saddle-core/src/caps.rs", "verified": true, "summary": "caps real; naming confuses."}, + {"title": "Load-time refuse lowered gemma4", "path_line": "carrier.rs + ar.rs", "verified": true, "summary": "Only generate-time error."} + ], + "changes": [ + {"title": "Fail-closed or wire lowered generate", "path_line": "carrier.rs + ar.rs", "cost": "hours", "summary": "Stop load-OK/run-fail."}, + {"title": "Optional bias Option/skip", "path_line": "weight_backend", "cost": "hours-days", "summary": "Unblock 31B-class."}, + {"title": "Fence/delete EAGLE product", "path_line": "speculative.rs", "cost": "days", "summary": "Reduce dual path."}, + {"title": "Move glimmer Carrier into arch crate", "path_line": "carriers.rs", "cost": "days", "summary": "Match maple/dots."}, + {"title": "Wire or delete glimmer CB", "path_line": "batch.rs + continuous_batch_route", "cost": "days-week+", "summary": "End half-migration."}, + {"title": "dots doc+vision+EOS", "path_line": "lib.rs spec_impl generate", "cost": "hours", "summary": "Cutover stale contract."}, + {"title": "Maple generate EOS", "path_line": "dense.rs maple", "cost": "hours", "summary": "Stop-token without Architecture."}, + {"title": "Retire toy as sole template", "path_line": "hipfire-arch-toy", "cost": "hours", "summary": "Point at maple+dots."} + ], + "report": "local://audit-ArchOther.md" +} +``` diff --git a/docs/plans/audit-2026-09-02/audit-ArchQwen.md b/docs/plans/audit-2026-09-02/audit-ArchQwen.md new file mode 100644 index 000000000..752e3f20f --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-ArchQwen.md @@ -0,0 +1,40 @@ + + +# Audit: ArchQwen + +slice=ArchQwen + +JSON contract: +{ + "slice": "ArchQwen", + "broken": [ + {"title": "Arch-6 EP MoE loads; generate routes dense TP", "path_line": "crates/hipfire-generate/src/qwen.rs:240", "verified": true, "summary": "5|6 → ep_serve_qwen35_dense_tp; dense_tp refuses num_experts!=0 (config.rs:214) and MoE layers (forward.rs:4168); forward_ep exists but unwired. #683."}, + {"title": "VL weight fail still sets vision_config", "path_line": "crates/hipfire-loader/src/carriers.rs:533-542", "verified": true, "summary": "load_weights map_err eprintln .ok(); has_vision_encoder only checks vision_config (lib.rs:1177) → image gate open with None weights."}, + {"title": "Multi-GPU DeltaNet EF residual empty", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/weights.rs:1746-1749", "verified": true, "summary": "new_with_quant_multi leaves s_ef_residual empty → stochastic DN; single-GPU Q8 EF default-on."}, + {"title": "DecodeBatchState hardcodes Q8 KV/DN", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/batch.rs:538-545", "verified": true, "summary": "new always Q8 filtered KV + StateQuant::Q8; reset_lane uses q8_lane_view only."}, + {"title": "dtype_from_quant_type subset of load quants", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/weights.rs:186-221", "verified": true, "summary": "Graded EP map misses load-admitted qts; fails closed as unsupported quant_type."}, + {"title": "Dense TP has no mrope/VL path", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/forward.rs:4132-4143", "verified": true, "summary": "forward_scratch_dense_tp lacks MropeCtx; VL single-device uses forward_scratch_mrope."}, + {"title": "forward_ep embed format allow-list narrow", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/ep_batch.rs:2180-2196", "verified": false, "summary": "Only HFQ4G256/128, Q8_0, F32; suspicious vs wider single-GPU embed."} + ], + "missing": [ + {"title": "ep_serve_qwen35_moe / admission refuse for arch 6 EP", "path_line": "crates/hipfire-generate/src/qwen.rs:240", "verified": true, "summary": "Arch implements forward_ep + EP batch state; generate never drives them. #683."}, + {"title": "Atomic VL config+weights publish", "path_line": "crates/hipfire-loader/src/lib.rs:2102-2103", "verified": true, "summary": "finish_qwen35_load assigns pair without cross-check."}, + {"title": "Multi-GPU EF + batch non-Q8 DN/KV", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/weights.rs:1746", "verified": true, "summary": "No multi EF; batch path not parameterized on StateQuant/KvMode."}, + {"title": "Serve/SlotEngine continuous-batch and EP bridge", "path_line": "crates/hipfire-arch-qwen35/src/serve_engine.rs:42-43", "verified": true, "summary": "Single-weight single-Gpu alternate path; CB deferred; no EP."}, + {"title": "Q4 DN tree-verify preflight", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/prefill.rs:4531-4536", "verified": true, "summary": "Runtime refuse only; no load-time gate."}, + {"title": "qwen2/llama shared transformer load extraction", "path_line": "crates/hipfire-arch-qwen2/src/qwen2.rs:22-28", "verified": true, "summary": "Duplicated load helpers; baseline capability deltas vs qwen35."}, + {"title": "Vision GPU quant + non-gfx1100 validation", "path_line": "crates/hipfire-arch-qwen35-vl/src/qwen35_vl.rs:351-370", "verified": true, "summary": "Host dequant HFQ4→F16; warn-only other archs."}, + {"title": "EP beyond frozen 4×gfx1201", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/ep_batch.rs:185-220", "verified": true, "summary": "rank_count!=4, REAP, paged refused."} + ], + "changes": [ + {"title": "Refuse arch-6 EP at admit or wire moe serve", "path_line": "crates/hipfire-generate/src/qwen.rs:240", "cost": "hours (refuse) / days–week (wire)", "summary": "Split 5 vs 6; drive forward_ep or fail before 19GB alloc."}, + {"title": "VL load fail-closed", "path_line": "crates/hipfire-loader/src/carriers.rs:533-542", "cost": "hours", "summary": "Pair config with weights; fix has_vision_encoder."}, + {"title": "Quant matrix test load↔dtype_from_quant_type", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/weights.rs:186", "cost": "hours–1 day", "summary": "Table every qt; document intentional EP unsupported set."}, + {"title": "Multi-GPU DN EF or explicit degrade", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/weights.rs:1746", "cost": "days (wire) / hours (degrade)", "summary": "Per-device residual + reset, or loud EF off on multi."}, + {"title": "Batch quant param or refuse non-Q8", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/batch.rs:538", "cost": "days / hours", "summary": "Align DecodeBatchState with carrier modes."}, + {"title": "Dense TP VL refuse or mrope", "path_line": "crates/hipfire-arch-qwen35/src/qwen35/forward.rs:4132", "cost": "hours / days", "summary": "Admit-time refuse cheaper than TP mrope."}, + {"title": "Document SlotEngine vs EP vs CB", "path_line": "crates/hipfire-arch-qwen35/src/serve_engine.rs:1-15", "cost": "hours", "summary": "Maintainer map of three multi-request shapes."}, + {"title": "qwen2/llama extraction + bias docs", "path_line": "crates/hipfire-arch-qwen2/src/qwen2.rs:22", "cost": "days+ / hours", "summary": "Shared transformer load; Dir bias policy clarity."} + ], + "report": "(full markdown in summary field; parent persists — scouts have no Write/local://)" +} diff --git a/docs/plans/audit-2026-09-02/audit-CliTuiScripts.md b/docs/plans/audit-2026-09-02/audit-CliTuiScripts.md new file mode 100644 index 000000000..fbcd0e6f7 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-CliTuiScripts.md @@ -0,0 +1,134 @@ + + +# Audit: CliTuiScripts + +# Audit: CliTuiScripts + +Scope: `crates/hipfire-cli/src/**`, `crates/hipfire-tui/**`, `scripts/**`, `tools/**`, `.github/workflows/{ci,no-gpu-ci,registry}.yml`. Read-only on master `/home/kaden/ClaudeCode/warpfront/hipfire`. + +## Broken + +### 1. TUI `/stats` uptime key mismatch — verified + +Serve emits `uptime_sec`; TUI reads `uptime_s` and defaults missing numeric keys to `0`, so Live serve Uptime always shows ~0s while queue/requests still populate. + +- Serve: `crates/hipfire-cli/src/serve/http.rs:469-482` — `"uptime_sec": meta.started.elapsed().as_secs()`. +- TUI: `crates/hipfire-tui/src/hipfire/dashboard.rs:455-470` — `uptime_s: as_u64("uptime_s")` with `unwrap_or(0)`. +- Render: `crates/hipfire-tui/src/ui.rs:468` — `fmt_uptime(s.uptime_s)`. +- How known: cross-read both sides. Valid object ⇒ `parse_stats` returns `Some` (not "stats unavailable"); only uptime is wrong. + +### 2. `stop --force` / `restart` system-wide `pkill -x daemon` — verified + +- `crates/hipfire-cli/src/serve/mod.rs:1399-1413`: `pkill -x daemon` (any exact-name daemon), optional `pkill -f target/release/hipfire-quantize`, `fuser -k {port}/tcp`. +- `crates/hipfire-cli/src/main.rs:660-678`: restart always `StopArgs { force: true }`. +- Docs warn (`docs/SERVE.md:47-58`) but restart defaults to force. Status discarded (`let _ = ...status()`). +- How known: read stop + restart + SERVE.md. + +### 3. Windows stop/ps/pid ownership is `/proc`-only — verified + +- `proc_start_time` `serve/mod.rs:1248-1251`; `pid_owns_listen_port` `:1258-1282`; `validate_serve_pid` `:1305-1316`; stop wait `:1385`; `kill -TERM` `:1376`. +- `ps` + aux: `main.rs:3437-3503` via `/proc`. +- On Windows plain stop: validate fails → deletes `serve.pid` without signaling → orphan serve. Force needs pkill/fuser. Additive to open #669/#643/#655. + +### 4. setup + launch_tui extensionless binaries only — verified + +- `setup.rs:285-305` installs `daemon`/`hipfire`/`hipfire-tui` without `.exe`. +- `launch_tui` `main.rs:3402-3414` probes `bin/hipfire-tui` only. +- Contrast: `find_daemon` `:5765-5769` prefers `daemon.exe`; `install.ps1` installs `*.exe`. +- Linux setup backup/rollback ordering is solid (stage→backup live→rename; reverse rollback; cleanup after install.json) — not broken. + +### 5. Dual automatic CI docs contradiction — verified + +| Source | Claim | +|---|---| +| `docs/VALIDATION.md:42-51` | Automatic = only `no-gpu-ci.yml` → `no-gpu-ci.sh` | +| `ci.yml:49-102` | Also on PR: workspace build, `cargo test --lib --workspace`, leanup-ratchets, crate-maps, ratchet-diff, cargo-deny | +| `CLAUDE.md:94-97` | Describes `ci.yml` | +| CONTRIBUTING + PR template | Push `no-gpu-ci.sh` | + +`no-gpu-ci.sh` tests a crate subset + change_gate/autoresearch Python; `ci.yml` is full workspace. Both fire on PRs. `registry.yml` is daily `registry_gen.py` only. + +### 6. PR template claims coherence-gate gone; still in tree — verified + +- PR template `:39-42`: "no longer exist in-tree". +- Present: `coherence-gate-{minimax,ornith15,qwen3-dspark,qwen35-dspark,cohere2moe,deepseek4-mtp,deepseek4-recall}.sh` + `_coherence_runner.py`. +- `gates.sh:7-9` correctly declines to call them. VALIDATION retired section accurate; PR template claim false. + +### 7. `bench --concurrency` batch backend always fails — verified + +- `bench_concurrency.rs:404-414`: `DaemonDriver::start` always `bail!`s (no multi-inflight Engine API). CLI still admits `--backend batch|both`. + +## Missing + +### 1. change_gate still wired while slated for deletion — verified + +- `no-gpu-ci.sh:25` unittest `tools.change_gate.tests.*` +- PR template `:24-34` required plan/run paste +- `leanup-thresholds.txt:33`, `docs/governance/2026-08-16-phase3-scope.md:75-89` +- Package: `tools/change_gate/**` +- VALIDATION does not list change_gate as automatic. Deletion without CI/PR-template cutover will red no-gpu-ci. + +### 2. hw-gate / hardware-evidence CI absent on master — verified + +- No `scripts/hw-gate/` (PR #679 path). `ci.yml:9-14` says GPU workflows removed. + +### 3. Large script corpus unreachable from CI/docs — structural + +- CI reaches ~8 scripts. VALIDATION/gates.sh maintain a small set. +- Orphan/historical clusters: coherence-gate-*, mi300x_*, dflash_diag_*, mtp_train/, reap/, campaign one-offs. No inventory of maintained vs historical. +- verified:false for per-file reachability; structural from entrypoint census. + +### 4. Windows lifecycle beyond install.ps1 — verified + +- install.ps1 parallel path (no setup atomic rollback/install.json writer). +- update Linux-only (`CLI.md:189`). +- stop/ps/restart Linux tooling (finding 3). + +### 5. autoresearch/ar/review still in no-gpu pytest — verified + +- `no-gpu-ci.sh:17-20` pytest `autoresearch/ar/tests`. Full `ar/review/**` remains. Same cutover need as change_gate if deletion planned. + +## Would change + +1. **hours** — Align `/stats` uptime keys (`http.rs:475`, `dashboard.rs:465`) + fixture test. Highest user-visible TUI bug. +2. **days** — Scope stop --force to hipfire-owned PIDs (`serve/mod.rs:1399-1413`); restart plain-stop first. +3. **days** — Windows process control + `.exe` install parity (setup, stop, launch_tui, ps). Coordinates with #669/#643/#655. +4. **days** — Unify VALIDATION/CLAUDE/PR-template with real CI; atomic cut of change_gate + fix coherence-gate in-tree claim; quarantine `coherence-gate-*.sh` under `scripts/historical/`. +5. **hours** — Refuse or HTTP-route `bench --backend batch` (`bench_concurrency.rs:404-414`). + +## Confidence + +Did: CLI Commands/setup/serve stop+http, TUI dashboard/ui/serve_ctrl, three workflows + no-gpu-ci/gates/install.ps1, VALIDATION/CLI/SERVE/PR template/CLAUDE, change_gate entry, coherence-gate samples, bench_concurrency dead path, script inventory vs CI refs. + +Did not: run tests/GPU; full 330-script callgraph; open gh issue bodies; full serve/complete.rs; every TUI chat SSE edge; deep tools/redline beyond CI unittests; autoresearch gate product semantics. + +## Parent JSON contract + +```json +{ + "slice": "CliTuiScripts", + "broken": [ + {"title": "TUI /stats uptime key mismatch", "path_line": "crates/hipfire-tui/src/hipfire/dashboard.rs:460-470; crates/hipfire-cli/src/serve/http.rs:469-482", "verified": true, "summary": "Serve emits uptime_sec; TUI reads uptime_s→0."}, + {"title": "stop --force / restart pkill -x daemon", "path_line": "crates/hipfire-cli/src/serve/mod.rs:1399-1413; main.rs:660-678", "verified": true, "summary": "System-wide pkill -x daemon + fuser -k; restart always force."}, + {"title": "Windows stop/ps is /proc-only", "path_line": "serve/mod.rs:1248-1385; main.rs:3437-3503", "verified": true, "summary": "Plain stop deletes pidfile without signaling on Windows."}, + {"title": "setup + launch_tui extensionless only", "path_line": "setup.rs:285-305; main.rs:3402-3414", "verified": true, "summary": "No .exe; diverges from install.ps1 and find_daemon."}, + {"title": "Dual automatic CI docs contradiction", "path_line": "docs/VALIDATION.md:42; ci.yml; no-gpu-ci.yml; CLAUDE.md:94", "verified": true, "summary": "VALIDATION names only no-gpu-ci; ci.yml also gates every PR."}, + {"title": "PR template says coherence-gate gone", "path_line": "PULL_REQUEST_TEMPLATE.md:39-42; scripts/coherence-gate-*.sh", "verified": true, "summary": "≥7 coherence-gate scripts still in tree."}, + {"title": "bench batch backend always bails", "path_line": "bench_concurrency.rs:404-414", "verified": true, "summary": "CLI admits batch/both; DaemonDriver::start always errors."} + ], + "missing": [ + {"title": "change_gate still wired", "path_line": "no-gpu-ci.sh:25; PULL_REQUEST_TEMPLATE.md:24-34; tools/change_gate/**", "verified": true, "summary": "CI+PR template require package slated for deletion."}, + {"title": "hw-gate absent on master", "path_line": "scripts/ (no hw-gate/); ci.yml:9-14", "verified": true, "summary": "PR #679 path not present."}, + {"title": "Script orphans", "path_line": "scripts/* vs workflows + VALIDATION", "verified": false, "summary": "CI ~8 scripts; large historical clusters."}, + {"title": "Windows lifecycle gaps", "path_line": "install.ps1; CLI.md:189; serve/mod.rs", "verified": true, "summary": "Parallel install; update Linux-only; stop Linux tooling."}, + {"title": "autoresearch/review in no-gpu pytest", "path_line": "no-gpu-ci.sh:17-20; autoresearch/ar/review/**", "verified": true, "summary": "Still load-bearing for CPU CI."} + ], + "changes": [ + {"title": "Align /stats uptime keys", "path_line": "http.rs:475; dashboard.rs:465", "cost": "hours", "summary": "Fix key + fixture test."}, + {"title": "Scope stop --force reaping", "path_line": "serve/mod.rs:1399-1413", "cost": "days", "summary": "Replace pkill -x; restart not always force."}, + {"title": "Windows pid/kill/.exe parity", "path_line": "setup.rs; serve/mod.rs; main.rs", "cost": "days", "summary": "taskkill + .exe names."}, + {"title": "Unify CI docs; cut change_gate wiring", "path_line": "VALIDATION.md; PR template; no-gpu-ci.sh", "cost": "days", "summary": "One merge-bar story; atomic deletion."}, + {"title": "Refuse or wire bench batch", "path_line": "bench_concurrency.rs:404-414", "cost": "hours", "summary": "Do not advertise dead arm."} + ] +} +``` diff --git a/docs/plans/audit-2026-09-02/audit-ConfigTopology.md b/docs/plans/audit-2026-09-02/audit-ConfigTopology.md new file mode 100644 index 000000000..2b59f406c --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-ConfigTopology.md @@ -0,0 +1,131 @@ + + +# Audit: ConfigTopology + +# Audit: ConfigTopology + +Checkout: `/home/kaden/ClaudeCode/warpfront/hipfire` @ `origin/master` `8cd15a62b` (read-only). + +Scope: `hipfire-config`; runtime `config.rs` / `loader_api.rs` / `multi_gpu.rs` / `device_mesh.rs` / `ep.rs`; `hipfire-registry` + `registry/`; CLI config/registry/pull/serve/bench. + +## Architecture (orientation) + +Ladder: `resolve(layers)` lowest→highest. Docs operator precedence: CLI one-shot > env > per-model > global TOML > registry > built-in. `load_env_layer` maps schema `env_compat` into typed keys; unknown `HIPFIRE_*` → `developer.` (`hipfire-config/src/lib.rs:3911-3953`). `ProcessConfig::from_resolved` keeps env_compat fields + developer keys; `legacy_value` / `developer_var` / `process_value` read the freeze only (`:3080-3088`, `:3224-3236`). `RuntimeConfig` OnceLock lowers once; `devices` rewritten to logical `0..N-1`. Device visibility: `hardware.devices`/`HIPFIRE_DEVICES` wins; else HIP/ROCR pair normalized (ROCR physical, HIP logical); `apply_device_visibility` `set_var`s both (`:3135-3145`). Multi-GPU: `Gpus` owns devices + `DeviceMesh`. `init_uniform`/`init_layers` → `from_parts` → mesh `Pp:N`; `init_tp` → mesh `Tp:N`; `single` → empty axes. EP loads all call `init_tp` (mesh labeled Tp, not Ep). `size_of(missing)=1`. Registry: v1 + curated `models.json`; dangling aliases silently retained-away; pull verifies sha/size when present. + +## Broken + +### 1. Production ambient HIPFIRE_* reads bypass ProcessConfig freeze (verified) +**path_line:** `docs/env-vars.md:51-69`; `crates/hipfire-daemon/src/main.rs:497-499,1199-1202,1575-1589,1801-1804,3294`; `crates/hipfire-loader/src/carriers.rs:409-413,2574-2578`; `crates/hipfire-runtime/src/calibration.rs:44-48,132-133`; `scripts/check-env-docs.py:31-76`; `scripts/no-gpu-ci.sh:29-30`. +**How known:** Docs claim production only reads bootstrap env. Grepped production `std::env::var("HIPFIRE_…")`. Checker fails any non-bootstrap literal outside CENTRAL_CONFIG_READERS/examples/tests. +**Impact:** Installed ProcessConfig can disagree with live ambient after daemon start. PP experimental compose and page-cache policy break the freeze story. Checker success string is false for this tree (CI status not re-executed). + +### 2. EP loads advertise Tp mesh; size_of(Ep)==1 (verified) +**path_line:** `crates/hipfire-loader/src/lib.rs:2844-2847,2993-2994,3221-3222,3360-3361,3491-3492`; `crates/hipfire-runtime/src/multi_gpu.rs:265-300`; `crates/hipfire-runtime/src/device_mesh.rs:192-200`. +**How known:** Every `load_model_ep*` uses `Gpus::init_tp`; init_tp sets `mesh = DeviceMesh::rect([(Tp, tp_size)])`; no Ep constructor; size_of defaults missing to 1. +**Impact:** EP still shards via `devices.len()` today. Mesh-aware #681 consumers asking Ep get single-rank. Topology twin of loader-admits / wrong-shape bugs. + +### 3. SpecLoadCfg documents loader env fallback that does not exist (verified) +**path_line:** `crates/hipfire-runtime/src/loader_api.rs:92-95`; `crates/hipfire-loader/src/spec_build.rs:201-202`. +**How known:** Comment says HIPFIRE_NGRAM_DRAFT* fallback and env wins; spec_build says never ambient, `unwrap_or(false)`. Daemon fills from RuntimeConfig/params (~1672) so CLI path OK — contract text is false. + +### 4. Ornith curated sampling pins disagree with registry test (verified) +**path_line:** `scripts/test_registry_gen_ornith.py:77-88`; `registry/models.json:1466-1512`. +**How known:** Test exact-equality without reasoning_effort; curated has xhigh on recommended/general/coding and none on instruct. Tag/files/aliases look correct; sampling pin drifted. + +### 5. Gpus::single vs from_parts(n=1) mesh shapes disagree (verified) +**path_line:** `crates/hipfire-runtime/src/multi_gpu.rs:238-262` (empty axes); `:828-847` (always Pp:n_devices). +**Impact:** Low today (size_of(Pp)=1; has_axis requires size>1). axes() length 0 vs 1 still diverges. + +### 6. CLI bench continuous_batch process-env side channel (verified) +**path_line:** `crates/hipfire-cli/src/main.rs:3925-3928,4007-4010`; schema `serve.continuous_batch_size` `hipfire-config/src/lib.rs:1013-1018`. +**Impact:** set_var/get bypasses ProcessConfig; checker-visible; schema already owns the knob. + +### 7. Experimental PP gates snapshotted then re-read live (verified) +**path_line:** `load_env_layer` maps HIPFIRE_PP_* → developer.*; daemon still `std::env::var` at `:1575-1589`. +**Impact:** Freeze contract broken for multi-GPU experimental compose. + +### 8. check-env-docs.py cannot truthfully pass on this tree (verified logic) +**path_line:** `scripts/check-env-docs.py:14-76` + Broken #1 sites. +Also misses non-literal reads, HIP_VISIBLE_DEVICES/ROCR_*, docs under docs/ not in REFERENCE_DOCS. + +## Missing + +### 1. No Gpus::init_ep / DimKind::Ep on EP admission (verified) +DeviceMesh supports Ep; constructors only emit Pp/Tp. EP overloads init_tp. Incomplete #681 cutover. + +### 2. hipfire registry undocumented in docs/CLI.md (verified) +CLI RegistryAction Status/List/Show/Update/Verify (`main.rs:303-325,642-643`). CLI.md config tables only; grep registry empty. + +### 3. init_vram_weighted stub (verified, documented) +`multi_gpu.rs:228-235`; `docs/multi-gpu.md:62`. + +### 4. Dangling registry aliases dropped silently (verified) +`hipfire-registry/src/lib.rs:328-336` retain without warning. + +### 5. Windows device-visibility hardening/docs (partial) +Unix-shaped set_var only (`:3135-3145`). Client child-env path better. No in-tree #669 string. Not hardware-proven broken. + +### 6. ModelEntry.sampling dual/inert vs recommended_settings (verified) +`lib.rs:196-199`; docs say sampling inert. Half-migration dead write surface. + +### Notes (not broken) +- No `parallel_capability.rs` on master loader tree — dual path is Tp-labeled EP, not two frameworks. +- Bare `hipfire config` → TUI (`main.rs:686-687`); docs/CLI.md:82 matches. `config ` defaults List (`:1076`). +- `hardware.devices` nullable string / HIPFIRE_DEVICES → ROCR physical + HIP logical — aligned. +- `HIPFIRE_PP_LAYERS` **is** wired via `developer_var` in carriers.rs (not dead). + +## Would change (ranked, with cost) + +1. **Route residual production env through ProcessConfig/developer_var/typed keys** (daemon PP/DFlash/DPM/log; loader PAGE_EVICTION; calibration). Honest bootstrap list in env-vars.md. Green check-env-docs = true. **Cost: 1–2 days.** +2. **Gpus::init_ep (or mesh-kind on init_tp) on all load_model_ep*** — DimKind::Ep; unit-test size_of(Ep). **Cost: hours–1 day.** +3. **Fix SpecLoadCfg/ngram contract** (+ optional RuntimeConfig default). **Cost: hours.** +4. **Reconcile ornith test with reasoning_effort** (separate instruct pin). **Cost: hours.** +5. **Delete bench HIPFIRE_BENCH_CONTINUOUS_BATCH bridge** — pass param. **Cost: hours.** +6. **Document hipfire registry in CLI.md; warn on dropped aliases.** **Cost: hours.** +7. **Unify Gpus::single mesh with from_parts(1).** **Cost: hours.** +8. **Windows visibility docs + prefer child-env injection.** **Cost: hours** docs; **days+** hardware matrix. + +## Confidence + +**Did:** config resolve/env/ProcessConfig/visibility; multi_gpu+device_mesh; EP loads; SpecLoadCfg vs spec_build; daemon/loader ambient reads; CLI config/registry/bench; registry parse+ornith; check-env-docs algorithm; docs env-vars/CLI/multi-gpu; production HIPFIRE greps. + +**Did not:** run checker/tests/builds; live HF sha/size; full gh novelty (hub capture incomplete); Windows hardware; deep ep.rs mesh consumers; every serve/run flag→ProcessConfig mapping; quant register beyond existence. + +**Novelty:** Did not confirm open-issue absence for env-bypass or EP-mesh labeling. #683 is generate-routing (peer). This EP finding is post-#681 **constructor labeling**, not the reverted device-mesh PR body. + +**Suspicious not verified:** whether no-gpu-ci is red on check-env-docs; whether generate already calls size_of(Ep); Windows #669 failure mode without issue body/hardware. + +## JSON summary (for parent synthesis) +```json +{ + "slice": "ConfigTopology", + "broken": [ + {"title": "Production ambient HIPFIRE_* bypass ProcessConfig", "path_line": "docs/env-vars.md:51-69; hipfire-daemon/src/main.rs:497,1199,1575-1589,1801; carriers.rs:409; calibration.rs:44; check-env-docs.py:31-76", "verified": true}, + {"title": "EP loads build Tp mesh; size_of(Ep)=1", "path_line": "hipfire-loader/src/lib.rs:2993,3221,3360; multi_gpu.rs:300; device_mesh.rs:192", "verified": true}, + {"title": "SpecLoadCfg env-fallback lie", "path_line": "loader_api.rs:92-95; spec_build.rs:201-202", "verified": true}, + {"title": "Ornith test vs curated reasoning_effort", "path_line": "test_registry_gen_ornith.py:77-88; registry/models.json:1482-1511", "verified": true}, + {"title": "Gpus::single vs from_parts mesh shape", "path_line": "multi_gpu.rs:238-262,847", "verified": true}, + {"title": "Bench continuous_batch set_var side channel", "path_line": "hipfire-cli/src/main.rs:3925,4007", "verified": true}, + {"title": "PP gates snapshotted then re-read live", "path_line": "hipfire-config load_env_layer:3911; daemon:1575-1589", "verified": true}, + {"title": "check-env-docs cannot truthfully pass", "path_line": "scripts/check-env-docs.py:57-76", "verified": true} + ], + "missing": [ + {"title": "No init_ep / Ep mesh admission", "path_line": "multi_gpu.rs:265-300; loader lib.rs:2844", "verified": true}, + {"title": "hipfire registry missing from CLI.md", "path_line": "cli main.rs:303-325; docs/CLI.md", "verified": true}, + {"title": "init_vram_weighted stub", "path_line": "multi_gpu.rs:228-235", "verified": true}, + {"title": "Silent dangling alias drop", "path_line": "hipfire-registry/src/lib.rs:328-336", "verified": true}, + {"title": "Windows visibility docs/hardening", "path_line": "hipfire-config/src/lib.rs:3135-3145", "verified": true}, + {"title": "Dual inert ModelEntry.sampling", "path_line": "hipfire-registry/src/lib.rs:196-199", "verified": true} + ], + "changes": [ + {"title": "Route residual env via ProcessConfig", "cost": "1-2 days", "path_line": "daemon/loader/runtime ambient sites"}, + {"title": "Gpus::init_ep for EP loads", "cost": "hours-1 day", "path_line": "multi_gpu.rs; loader EP"}, + {"title": "Fix SpecLoadCfg contract", "cost": "hours", "path_line": "loader_api.rs; spec_build.rs"}, + {"title": "Ornith reasoning_effort pins", "cost": "hours", "path_line": "test_registry_gen_ornith.py"}, + {"title": "Delete bench env side channel", "cost": "hours", "path_line": "cli main.rs:3925-4010"}, + {"title": "CLI.md registry + alias warnings", "cost": "hours", "path_line": "docs/CLI.md; registry parse"}, + {"title": "Unify single-device mesh", "cost": "hours", "path_line": "multi_gpu.rs:238-262"}, + {"title": "Windows visibility docs/child-env", "cost": "hours-days", "path_line": "config visibility; docs"} + ] +} +``` diff --git a/docs/plans/audit-2026-09-02/audit-Dflash.md b/docs/plans/audit-2026-09-02/audit-Dflash.md new file mode 100644 index 000000000..13c564f12 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-Dflash.md @@ -0,0 +1,59 @@ + + +# DFlash speculative-decode audit — 2026-09-03 + +_Lifecycle: planned intent. Read-only audit of master `8cd15a62b`. Findings cite `path:line` at that commit._ + +**Scope.** Five read-only slices: draft runtime + VRAM (`hipfire-runtime/src/dflash.rs`, `admission.rs`, the three `carrier.rs` scratch-sizing commits), the generate loop (`hipfire-generate/src/qwen.rs` `generate_dflash`/`generate_spec` + helpers), verify/replay (`hipfire-arch-qwen35/src/dflash_spec.rs`, `dflash_verify_pm4.rs`, `speculative.rs`), the `attention_dflash*` kernel family + dispatch, and pairing/discovery/config (`dflash_generic.rs`, `dflash_convert.rs`, daemon/CLI/config). The headline finding (#1) was re-derived by the auditor from `crates/hipfire-cli/src/main.rs:1944-1965,2488+` and `crates/hipfire-daemon/src/main.rs:1194-1216`. + +## Verdict + +The speculative machinery itself is in good shape: the verify block is a batched target prefill with a DeltaNet snapshot + innovation tape replayed for `accept_len+1` (not per-row checkpoints), graph capture goes through owned kernarg blobs (`launch_maybe_blob`), retained PM4 is a pure phase machine at fixed B=16, the draft's attention is F32 end to end and never touches the target's quantised KV, masks and online-softmax are correct in the live kernels, and the client-cancel path runs the canonical rollback. The problems are around it: **reachability, accounting, and a handful of exits that skip rollback.** + +## Broken + +1. **DFlash is unreachable without an explicit draft path — the documented on-ramp does not exist.** high, verified. + `AGENTS.md` ("pull the draft, `hipfire config set dflash_mode auto`, run; expect `[hipfire] DFlash draft detected`") and `docs/MODELS.md:116` describe sibling-filename auto-discovery (`qwen3{ver}-{size}-dflash-{quant}.hfq`). No such matcher exists in daemon, CLI, loader, or registry. The only draft sources are `--model-draft` (`cli/main.rs:1960-1965`), `developer.dflash_draft` / legacy `HIPFIRE_DFLASH_DRAFT` (`cli/main.rs:2579-2604`, daemon `main.rs:1194-1216`), and `params.draft`. `load_params` (`cli/main.rs:2488+`) never consults the registry entry, and `hipfire pull -draft` sets nothing. Consequence: `dflash_mode auto` (and `on`) with a pulled draft runs plain AR, silently — `on` behaves identically to `auto` on a missing/failed draft (daemon strips the path only for `off`; Qwen35 load `Err` → `eprintln` + AR, `loader/lib.rs:1960-1972`). The registry already carries the pairing (`qwen3.8:27b-draft` → `qwen3.8:27b-draft-mq4` → `qwen38-27b-dflash-mq4.hfq`), so a registry-driven resolution in `load_params` is the whole fix; `on` should then fail the load when no draft resolves, mirroring MTP (`lib.rs:1973-1976`). Measured stakes on a 7900 XTX: 27B mq4 AR is 46.8 tok/s; every "150 tok/s" number users quote is a speculator. +2. **Draft VRAM is charged nowhere.** high. `ModelFootprint` has two fields (`admission.rs:17-22`: target weights, KV bytes/token); nothing adds draft weights (0.92–1.66 GiB by quant, `docs/MODELS.md:106-113`) or the L-indexed draft planes. At the default `HIPFIRE_DFLASH_CTX_CAP=8192` on 27B (h=5120, ne=5, kvd≈1024, nL=5, B=16): `target_hidden` L·ne·h·4 ≈ 800 MiB, `target_hidden_proj` ≈ 160 MiB, per-layer K/V ctx caches ≈ 320 MiB, `k/v_cat` ≈ 160 MiB, `mq_x_rot` ≈ 100 MiB (chunked; was 800 MiB before the chunk fix) — **≈1.5–1.8 GiB scratch + weights ≈ 2.5–2.7 GiB** invisible to admission. Uncapped at 32k the caller's own comment puts it at ~11 GB (`dflash_spec.rs:117-122`). Windowed DFlash2 drafts (all-sliding, W=2048) are much smaller. Separately: the daemon never constructs `AdmissionController` at all — its only users are the slots path (`session_table.rs`) and tests — so `noslots` serving has no budget gate; the user gets a raw allocation failure after a 19 GB load. +3. **Partial-failure leaks in the draft constructors.** high. `DflashScratch::new_with_mq` (`dflash.rs:1403-1483`) and `new_windowed` (`1343-1356`) `?` out mid-way without freeing earlier `alloc_tensor`s; `DflashWeights::load` (`637-946`) likewise; `dflash_generic::build_generic_dflash_speculator` loads weights then `?` on scratch without freeing (`dflash_generic.rs:1048-1055`). `GpuTensor`/`DeviceBuffer` have no `Drop` (`rdna-compute/src/dispatch.rs:205-211`). The Qwen35 outer chain is transactional (`or_free!` in `load_dflash_state`), the ctors it calls are not. +4. **`make_spec_emitter` Err after a successful prefill has no rollback.** high. `qwen.rs:3113-3125` emits a validation error and returns `None`; the target's KV/DeltaNet/drafter hidden already advanced and, on `!cache_hit`, host `seq_pos`/`conversation_tokens` were cleared at `2929-2936`. Next turn can LCP against a dirty GPU. Every other error exit in the loop (`prefill` Err, `step` Err, realign, forced terminal, pending-seed flush, mid-loop abort) runs `production_fail_closed_rollback_live` — this one is the odd one out. +5. **Entry cap vs loop cap disagree.** high. `generate_dflash` falls through to AR only when `prompt + max_tokens > ctx_capacity` (`2078-2090`); `generate_spec` then hard-errors when `prompt + max_tokens + block_size > ctx` (`2980-3000`) — after `gen_start` was emitted. Requests in that `block_size` band get a started generation followed by an error instead of the promised AR fallback. Related: the mid-loop `position + block_size >= ctx_capacity` `break` (`3358-3360`) sets no flag, so if `generated < max_tokens` the epilogue reports `finish_reason=stop` and may store the cache — a silent early stop (med). +6. **Draft pairing checks nothing about the target.** high. Convert records draft geometry only (`dflash_convert.rs:1031-1088`: block size, mask token, `target_layer_ids`, `num_target_layers`, dims; no target family/size/quant/hash). Qwen35 load checks only `target_layer_ids[i] < n_layers` (`dflash_spec.rs:333-344`) — not even `draft.hidden == target.dim` (llama's generic path does, `carriers.rs:944-951`); `num_target_layers` is written and never compared. The known 3.5-draft-on-3.6-target τ≈1.2 failure is this. Also `vocab_size` is written as `config.get("vocab_size").cloned()` (`dflash_convert.rs:1040`) and can be JSON `null`, which `DflashConfig::from_hfq` (`dflash.rs:134`) then refuses — the writer can emit a draft the loader cannot read. + +## Missing + +- Registry-driven draft resolution at load (fixes 1); `dflash_mode=on` failing closed on a missing/failed draft. +- `ModelFootprint` (or a side charge at load) covering draft weights + `scratch(L)`; a "ceiling at this config is N tokens" line before the load; an admission test that `admit(64k)` with 27B + draft on 24 GiB fails closed (fixes 2). +- Transactional ctors (mirror `or_free!`) in `dflash.rs`/`dflash_generic.rs` (fixes 3). +- One shared predicate for the entry cap and the loop cap; a `ctx_exhausted` stop reason for the mid-loop break that suppresses cache store (fixes 5). +- Target identity in the draft header (arch, n_layers, dim, vocab, optional weight hash) + a loader refusal; `num_target_layers == target.n_layers` and `draft.hidden == target.dim` on the Qwen35 path; non-null `vocab_size` in convert; a schema version for the dflash metadata beyond `HFQ_VERSION=1` (fixes 6). +- `GenericDflashSpeculator` inherits no-op `rewind_to`/`on_evict` (`spec.rs:727-729, 807-818`; `dflash_generic.rs:951-968`) — prompt-cache resume / CASK on the llama path can desync `target_hidden_host` from the target KV. Either implement or refuse those combinations. +- Kernels: `attention_dflash.hip:254-259` finalizes with `1/l_run` and no `l_run == 0` guard (the sliding sibling guards, `:285-297`); the `n128_f16kv` entry is registered under `HasWmma` (includes gfx12) while its source is gfx11-only (`attention_table.rs:370-385`) — shadowed today by the higher-priority `DflashV5Gfx12`, a JIT failure if that ordering changes; the production `attention_dflash_wmma*` launches pass stack `void**` kernargs (`attention.rs:8588-8616, 9306-9336, 9440-9454`) — safe only because the draft captures FFN, never attention. + +## Would-change + +- `dflash_adaptive_b` is parsed by the daemon and dropped (`let _adaptive_b`, `daemon/main.rs`); the config key (`speculation.dflash_adaptive_b`, default true) is inert. Adaptive B lives only in `examples/dflash_spec_demo.rs`. Either wire or delete the key. +- AR fallback past the ctx cap leaves the full draft allocation resident (`qwen.rs:2078-2090`). Reasonable for speed; wrong on a 24 GB card after a load that only just fit. +- Windowed split drafts still pin one full-reach K/V layer at `w_full = requested_ctx` (`dflash_spec.rs:294-304`, `dflash.rs:1349-1352`) — the same "logical max_seq" class the carrier fix retired; consider `physical_cap`. +- `attention_dflash_wmma_m32_kstg_FAILED.hip` is unwired but exports the same C symbol as the live `m32` kernel; an accidental `include_str!` would shadow it. +- The "identical output" claim on AR fallback (`qwen.rs:2083`) is true for the main Qwen route (same request seed); the `Qwen2Spec` fallthrough uses a bare `generate_qwen2` without tools/seed parity (`ar.rs:1348-1372`). +- Docs: `AGENTS.md` §1/§3.6/§7 and `docs/MODELS.md:116` describe the non-existent auto-match; `docs/env-vars.md` should list `developer.dflash_draft` as the only production knob until (1) lands. + +## Confirmed by design + +- Verify forward: DeltaNet advances in place, then restores the pre-window snapshot and replays the `GdnTape` innovations for `accept_len+1` (`speculative.rs`); no per-row S checkpoint is needed. +- Graph capture: verify graphs are cached per B via `capture_mode → launch_maybe_blob` with owned blobs (`dispatch.rs`, `graph.rs`); the draft captures only its FFN tail with owned blobs (`dflash.rs:1969-1994`). Retained PM4 pins B=16, kv_mode/dn_state q8, `!tree`, `!full_logits` (`dflash_verify_pm4.rs`). +- Client cancel after handshake runs `production_fail_closed_rollback` (`qwen.rs:2557-2561, 2724-2727`) — not a bare `Abort => {}`. +- The three scratch-sizing commits end with the target's flash partials sized from `kv.physical_cap` (`carrier.rs:85-88`); `c305a34b6` was a bad intermediate that `a22b88d3f` restored. No sibling in `dflash.rs` is sized the old way. + +## Not read + +Full `draft_forward` attention body (~`dflash.rs:2500-3100`), the MTP speculator's commit path, the vision DFlash loop, research-only attention variants (`n64`, `v2/v3-noncausal/v4/v6/v7/v7b`) beyond status, `dflash_spec_demo.rs`, and the numerical tests under `rdna-compute/examples/test_attention_dflash.rs`. Nothing was run. + +## Recommendation, in order + +1. Registry-driven draft resolution + `on` fails closed + docs (hours; the largest measured user-facing win available: AR → speculative for every default user). +2. Charge the draft in the footprint and print the context ceiling at load (hours); decide whether `noslots` serving should construct `AdmissionController` at all (design). +3. Transactional draft ctors; `make_spec_emitter` rollback; unify the two ctx-cap predicates and add the `ctx_exhausted` reason (hours each, one PR). +4. Target identity in the draft header + loader refusal; non-null `vocab_size` (a day; needs a convert re-run for existing drafts or a permissive read of old headers). +5. Kernel nits (zero-`l` guard, `n128` predicate, blob-safe attention launches) in one small PR. diff --git a/docs/plans/audit-2026-09-02/audit-DispatchKernels.md b/docs/plans/audit-2026-09-02/audit-DispatchKernels.md new file mode 100644 index 000000000..33bbd9740 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-DispatchKernels.md @@ -0,0 +1,99 @@ + + +# Audit: DispatchKernels + +## Broken + +### 1. MQ8 family path: rotate does not produce INT8 scratch; prerotated unwraps (VERIFIED) + +**Citations:** `crates/hipfire-dispatch/src/types.rs:130` (MQ8G256 => RotationPlan::Mq8Internal); `crates/hipfire-dispatch/src/families/gemv.rs:143-148` (Mq8Internal o WithRmsnorm or Plain); `crates/hipfire-dispatch/src/families/rotation.rs:62-199` (no MQ8 arm; Plain o rotate_x_mq F32 FWHT); `crates/hipfire-dispatch/src/families/gemv.rs:246-260` (run_input Raw o rotate then post-rotation Prerotated); `crates/hipfire-dispatch/src/families/gemv.rs:512` (gemv_mq8g256_prerotated(w.buf, y, m, k) — no x); `crates/rdna-compute/src/gemv.rs:6151-6164` (docs require prior rotate_quantize_x_mq8; mq_x_q8.as_ref().unwrap()). + +**How known:** Cross-read rotation plan o select_rotation_variant o RotationFamily::run o launch. Pipeline path is fixed: `crates/hipfire-dispatch/src/pipeline/steps.rs:940-948` special-cases Mq8Internal to rmsnorm_f32 + rotate_quantize_x_mq8. Correct low-level API is gemv_mq8g256_with_rotate (gemv.rs:6178-6204). Family run_auto/rotate never calls it. + +**Effect:** GemvFamily::run_auto / RotInput::Raw on MQ8 panics on unwrap. Tests leave RotationPlan::Mq8Internal empty (tests.rs:770). + +### 2. Plain KernelKey::GemvMq8G256 (and other plain MQ keys) missing from launch (VERIFIED) + +**Citations:** types.rs:703 maps (MQ8G256, Plain) => GemvMq8G256; gemv_table.rs registers MQ8 plain+prerotated; gemv.rs:537 other => MissingImpl. + +**How known:** Exhaustive match inspection — only GemvMq8G256Prerotated handled. run_auto uses post-rotation Prerotated so hits bug #1; direct Plain o MissingImpl. + +### 3. gemm_table V2 comment contradicts registrations and sources (VERIFIED contradiction) + +**Citations:** gemm_table.rs:364-370 claims MQ6/5/3/2V2 and MQ4CG256 remain gfx12-only while registering GemmMq{6,5,4,3,2}G256V2* as ArchPredicate::HasWmma (368-482). MQ4C correctly uses HasWmmaGfx12 (489-506). kernels.rs:3124-3141 includes gfx11 residual WMMA sources for MQ4/5/6/3/2 V2. Runtime gemm_mq6g256v2_residual_wmma branches gfx12 then gfx11 (gemm.rs:32190-32210). + +**How known:** Comment vs register vs include_str! vs residual dispatcher. Not a silent wrong route on WMMA GPUs; documentation/registry intent drift. + +### 4. MQ6 residual module name suffix _mq5v2 (VERIFIED smell) + +**Citations:** gemm.rs:32097 format!("{}_mq5v2", kname) inside gemm_mq6g256v2_residual_wmma_gfx12; same pattern at 31884, 33308, 34558. Contrast QKV MQ6: module_v2 = "…_mq6v2" (31639). + +**How known:** Grep _mq5v2. Functional collision avoided because kname embeds mq6 vs mq5. Copy-paste defect. + +### 5. verify-bind-thread.sh does not cover most impl Gpu surfaces (VERIFIED) + +**Citations:** Script default FILE=crates/rdna-compute/src/dispatch.rs (line 21); finds single impl Gpu { … impl Drop for Gpu (40-45). Actual impl Gpu also in gemv.rs:171, gemm.rs:168, attention.rs:164, moe.rs:13, norm.rs:90, embedding.rs:13, sampling.rs:95, gemma4_ext.rs, etc. + +**How known:** Grep ^impl Gpu + script read. Multi-GPU half only flags first call after `let = &mut gpus.devices[X]` if it is `.hip.` — misses other patterns. + +--- + +## Missing + +### 1. RotationVariant / RotationFamily support for MQ8 INT8 quantize (VERIFIED gap) +Root of broken #1. Only pipeline special-case implements the contract documented on rotate_quantize_x_mq8. + +### 2. Automated bind_thread proof for gemv/gemm/attention/moe (VERIFIED gap) +Spot checks show many self.bind_thread()? first lines in gemv.rs; script cannot fail CI on regressions outside dispatch.rs. + +### 3. DeviceBuffer RAII Drop (VERIFIED design gap) +hip-bridge malloc returns owning buffer; free is explicit (ffi.rs:911-924); no Drop for DeviceBuffer. SAFETY docs on from_raw/alias/from_vmm_owner present (lib.rs:114-147). Leak-on-forget, not double-free for borrowed. + +### 4. hsa-bridge per-site SAFETY comments (VERIFIED gap) +hsa-bridge/src/lib.rs dense unsafe for init, agents, queues, signals without // SAFETY: blocks. Experimental Phase-2 surface. + +### 5. MoE shape_gate: BatchGt(1) not enforced by resolve (documented intentional) +moe_table.rs:12-20,40-45. Prefill executor dtype/env dispatch bypasses registry. + +### 6. CK feature optional soft-fail at load; try_* fail-closed (mostly OK) +**Citations:** feature flash-attn-ck (rdna-compute lib.rs:14-15); load soft-fail WARN (dispatch.rs:1157-1172); reject matrix Decode/SmallQuery/NonCausal/Graph/Replay/Tree/Window/Block/CapabilityMiss (flash_attn_ck.rs:101-116,131+); attention family cfg + HIPFIRE_FLASH_PREFILL force-off (attention.rs:743-755). try_* returns Ok(false) o native path. **Verified fail-closed at launch selection**; load failure soft by design. + +### 7. Fused MQ V2 table Always vs GEMM WMMA gates (coverage clarity missing) +fused_qkv_table.rs:17-22 Always for FusedQkvMq*V2 and MQ4C. Fused impls load scalar HIP (fused_qkv_mq6g256v2 at gemm.rs:32323+). GEMM MQ4C Err without gfx12 (35461-35468). + +--- + +## Would change (ranked) + +1. **Wire MQ8 through family rotation + launch plain arm** — hours + Add RotationVariant::Mq8Quantize o gpu.rotate_quantize_x_mq8; fix select_rotation_variant; map plain key to gemv_mq8g256_with_rotate or delete plain registration; GPU-free test that rotate fills scratch; regression against steps.rs Mq8 path. + +2. **Extend verify-bind-thread.sh to all impl Gpu modules** — hours + Plus broader multi-GPU patterns. Highest leverage against multi-device silent wrong-device allocs. + +3. **Fix gemm_table V2 comment + _mq5v2 suffixes** — hours + Cheap correctness-of-docs and module-cache clarity. + +4. **Fused V2 Always contract test / or HasWmma gate** — hours + Prevent admit-then-JIT-fail if fused sources are arch-limited. + +5. **DeviceBuffer Drop or must_use discipline** — days + Coordinate VMM owners; reduce leak class. + +6. **CK first-miss diagnostics** — hours + Keep soft-fail load; make capability/workspace misses obvious once. + +7. **hsa-bridge SAFETY comment pass** — hours + Low urgency until HSA path is production-hot. + +--- + +## Confidence + +**Did:** Core dispatch tables (gemm/gemv/moe/fused_qkv), GemvFamily+RotationFamily+pipeline MQ8 path, gemm MQ V2 residual/plain/gfx11/gfx12, kernels.rs V2 includes, CK load+select+attention cfg, bind_thread script, hip-bridge DeviceBuffer ownership, hsa-bridge unsafe survey, fused MQ4C Always vs gemm gfx12. + +**Did not fully:** Exhaustive (op imes qt imes arch) matrix for every KernelKey; dead kernels/src/*.hip vs include_str! full set-diff; every impl Gpu bind_thread first-stmt proof; WMMA gfx11 vs gfx12 numerical drift inside HIP sources; open GitHub issue cross-check via gh (unavailable this turn — novelty not claimed against issue tracker); live GPU/JIT validation. + +**Suspicion (not verified broken):** gate_up HFQ6/MQ6 fusion requiring dp4a_eligible may leave non-dp4a arches on slower unfused path only — appears intentional. MQ6 GEMV on gfx906 uses HFQ6 generic path (gemv_mq6g256_prerotated o gemv_hfq6g256) — not proven wrong without arch kernel internals. + +**Architecture note:** hipfire-dispatch KernelRegistry+families resolve (op,dtype,arch) o rdna-compute Gpu methods (impl Gpu split across gemv/gemm/attention/moe/…) JIT via kernels.rs include_str! o hip-bridge launch; optional flash-attn-ck sidecar under feature; bind_thread intended on every device-touching Gpu entry; rotation plans (FwhtG256/Mq8Internal/Givens) sit above GEMV prerotated kernels. diff --git a/docs/plans/audit-2026-09-02/audit-DocsIssues.md b/docs/plans/audit-2026-09-02/audit-DocsIssues.md new file mode 100644 index 000000000..530a48ef5 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-DocsIssues.md @@ -0,0 +1,144 @@ + + +# Audit: DocsIssues + +# Audit DocsIssues FULL REPORT + +Slice DocsIssues. Checkout /home/kaden/ClaudeCode/warpfront/hipfire @ origin/master 8cd15a62b. gh issues/PRs 2026-09-02. No local:// written (no Write tool). + +## Broken + +1. Dead --example daemon/test_kernels (verified). CONTRIBUTING.md:58-59 GETTING_STARTED.md:129 VALIDATION.md:84-97 prescribe cargo build --release --features deltanet --example daemon -p hipfire-runtime and --example test_kernels. Product: crates/hipfire-daemon Cargo.toml:8-10 [[bin]] daemon. AGENTS/CLAUDE already correct. No test_kernels example/source on master (workspace grep empty). Impact: contributor dead ends; open #645 NixOS no example named. + +2. Skills path .skills/ missing (verified). CONTRIBUTING.md:33-44,139-145,269-273 → .skills/hipfire-*. Actual .agents/skills/ (INDEX+CLAUDE+disk). .skills/ absent. + +3. MODELS DeepSeek default wrong (verified). MODELS.md:148-149 bare deepseek-v4-flash → mq2lloyd 86.2GB. registry/models.json:255-257 default file mq2r 82GB; Lloyd is deepseek-v4-flash:mq2lloyd :214-216. MQ2R default since 2026-08-14 per registry desc. + +4. README counts disagree (verified). Badge 61 curated; body 77 pullable; models.json ~79 keys (maple-preview, bonsai bq1/tq2, ornith, Qwen3.8 ladder). README family table omits maple/bonsai/ornith. hipfire list -r is correct authority; marketing numbers false. + +5. ARCHITECTURE force_local wrong (verified). ARCHITECTURE.md:107: HIPFIRE_LOCAL|--kv-mode|--json|--no-stream. CLI.md:70: JSON/non-stream do NOT force local; HIPFIRE_LOCAL|--kv-mode|--image. Source main.rs:1911-1918: HIPFIRE_LOCAL|image|kv_mode|kv_backend|speculation|model_draft|draft_max|dspark_conf_threshold — neither json nor no_stream. CLI closer but incomplete. + +6. ARCHITECTURE carriers incomplete (verified). Table ends Cohere2Moe(12). Missing Gemma4(13/22), Muse Glimmer(14), Maple(15). architecture-ids + workspace crates present. + +7. INDEX ownership/pin drift (verified). INDEX.md:8-12 inventory date 2026-07-22 branch beta audited ref 202282de…. Release line master. Truth collision: INDEX perf-checkpoints=measured BENCHMARKS=historical; perf-checkpoints/README.md:3-9 lifecycle historical every file AND calls BENCHMARKS current product claims; BENCHMARKS.md:1-10 self-historical. Missing from INDEX top-level set while linked: crate-maps.md GLOSSARY.md GEMMA4_ESERIES_* qwen35-vl-mq4v2-spec.md lfm2-vl-mq4v2-spec.md. + +8. CONTRIBUTING active asks closed (verified). CONTRIBUTING.md:286-295 Issue #57 CLOSED 2026-04-27; #58 CLOSED 2026-05-06; #50 OPEN title Test on gfx1152 (not crash prose). Crate topology omits daemon/loader/engine/generate/config/registry/cli/most arch crates. + +9. Validation retired harness (verified). VALIDATION.md kernel channel still test_kernels; deepseek4-pr-body.md:87-90 historical coherence-gate+test_kernels easy to re-copy. ci.yml agentic GPU gates removed; GPU manual per VALIDATION. + +## Missing + +1. No docs drift CI beyond scripts/check-env-docs.py (VALIDATION). No link/example/INDEX completeness gate. +2. MODELS/README incomplete vs registry; no single Gemma status owner across #672/#667/#678/#614/#270. +3. Issue hygiene: 88 open; ~50+ stale >30d (#30–#560 band). #683 undoc EP generate non-route. #681 G1 DeviceMesh merged 2026-09-02; multi-gpu.md may lag. #648 overlaps #682 G2. +4. PR garden 14 open: #682 G2 MERGEABLE hot; #680 ornith alias; #679 hw-gate MERGEABLE rewrites VALIDATION/CONTRIBUTING (doc ownership conflict); #677/#675 draft CONFLICTING historical post-revert do-not-merge; #670 maple CONFLICTING large; #667 gemma draft CONFLICTING; #652 dflash draft stale fixes #640; #643 Windows VMM draft stale fixes #635; #595 UTF-8 draft base beta CONFLICTING stale touches old examples/daemon; #593 tool_calls draft base beta; #563 DS4 paging draft base beta CONFLICTING partial supersession; #527 mesh megabranch superseded forensic. Same-surface #682/#683: #682 #670 #667 #595 #563 #677 #675; #679 docs-only. + +## Would change (ranked, cost) + +1. hours — Fix dead entrypoints: CONTRIBUTING+GETTING_STARTED+VALIDATION → cargo build -p hipfire-daemon; drop/replace test_kernels or mark kernel channel blocked; .agents/skills/; replace asks with #666/#683/#669/#645. +2. hours — MODELS.md:148 mq2r default; README drop fake 61/77 or generate from models.json; add maple/bonsai/ornith or mark partial. +3. hours — force_local: CLI.md sole owner match main.rs:1911-1918; ARCHITECTURE link not invent --json/--no-stream. +4. hours-day — INDEX refresh master pin; list missing pages; checkpoints=measured evidence; BENCHMARKS=historical only; fix perf-checkpoints README. +5. day — ARCHITECTURE carriers Gemma/Muse/Maple; document #683 EP generate known non-route; sync #681 DeviceMesh if public. +6. day — Issue triage: close #155 announcement; stale label; collapse dups #537↔#646 multi-slot, #623↔#669 GPU, #162↔#645 NixOS, #270↔#672↔#614↔#678↔#667 Gemma, #640↔#652↔#459 DFlash, #448↔#563 pager, #50↔#353 gfx1152. +7. days — Docs drift CI: forbid dead --example and .skills/ links; INDEX completeness; optional models.json badge gen. +8. week+ — PR garden: close #527/#675/#677 historical; rebase/close beta #593/#595/#563; sequence #682 then mesh G3/G5 per #666; land #679 with VALIDATION ownership handoff. + +## Open issues table + +| # | Cluster | Age | Status | Recommendation | +|---|---|---|---|---| +| 683 | bug/EP-MoE generate | hours | open hot | P0 document+fix archetype | +| 678 | bug/gemma4 | ~1d | open | triage #614/#667/#672 | +| 672 | feat/gemma publish | ~1d | open | keep block artifact | +| 669 | bug/GPU select | ~1d | open | dup-check #623 | +| 666 | tracking/device-mesh | active | open | sole mesh authority post-revert | +| 655 | bug/Windows dump | ~1d | open | split actionable | +| 651 | feat/quant PARO | ~6d | open | backlog | +| 650 | research/KV | ~6d | open | research | +| 649 | research/redline | ~6d | open | research | +| 648 | bug/EP constructors | ~6d | open | align #682 G2 | +| 647 | roadmap/adaptive KV | ~6d | open | roadmap | +| 646 | feat/multi-slot serve | ~6d | open | merge #537 | +| 645 | bug/NixOS examples | ~7d | open | same CONTRIBUTING rot | +| 644 | bug/MQ4V2 decode | ~7d | open | P0 correctness | +| 642 | feat/IQ3_S | ~7d | open | backlog | +| 640 | bug/DFlash leak | ~6d | open | PR #652 | +| 639 | bug/Windows paths | ~9d | open | Windows cluster | +| 635 | bug/Windows VMM | ~7d | open | PR #643 | +| 623 | bug/Windows iGPU | ~11d | open | cluster #669 | +| 614 | bug/gemma hipGraph | ~10d | open | gemma cluster | +| 605 | RFC/discovery API | ~13d | open | RFC | +| 604 | RFC/response_format | ~13d | open | RFC | +| 588 | bug/pp=2 exactness | ~18d | open | multi-GPU | +| 587 | bug/redline GC12 | ~17d | open | needs-triage | +| 577 | bug/physical_cap | ~22d | open | long-ctx | +| 569 | bug/redline-rocr | ~27d | open | redline | +| 568 | bug/RefCell | ~27d | open | runtime | +| 560 | call/regression testers | ~1mo | open | community | +| 558 | feat/gfx1030 branch | ~1mo | open | stale? | +| 540 | feat/tokenizer | ~1mo | open | backlog | +| 537 | feat/multi-slot | ~27d | open | dup #646 | +| 533 | feat/MQ2 prefill | ~1mo | open | backlog | +| 526 | fix/CLI extensions | ~1mo | open | check fixed | +| 499 | bench/W7900 | ~1mo | open | accept or close | +| 491 | bug/instrument | ~1mo | open | tool | +| 490 | bug/instrument | ~2mo | open | stale | +| 486 | bug/DFlash OOM | ~27d | open | spec-decode | +| 478 | bug/gfx1030 default | ~1mo | open | RDNA2 | +| 475 | bench/gfx1201 | ~2mo | open | stale | +| 469 | feat/DFlash trainer | ~2mo | open | research | +| 462 | bug/DeltaNet serve | ~2mo | open | check post-revert | +| 461 | bug/llama dir dispatch | ~2mo | open | verify | +| 459 | bug/DFlash draft load | ~2mo | open | cluster #640 | +| 456 | tracking/MoE lattice | ~2mo | open | research | +| 448 | tracking/weight pager | ~2mo | open | #563 | +| 443 | proposal/AUR | ~2mo | open | community | +| 433 | question/status | ~2mo | open | answer+close | +| 392 | feat/AWQ+GPTQ | ~2mo | open | relabel if shipped | +| 354 | research/HBW-KV | ~3mo | open | stale | +| 353 | bench/gfx1152 | ~3mo | open | pairs #50 | +| 346 | feat/Orthus MTP | ~3mo | open | backlog | +| 345 | question/Nemotron | ~2mo | open | answer+close | +| 344 | research/FeatherOps | ~3mo | open | stale | +| 343 | roadmap/AWQ trunk | ~3mo | open | partial ship? | +| 341 | research/AWQ calib | ~3mo | open | research | +| 328 | refactor/rdna-compute | ~3mo | open | chore | +| 305 | governance/dual-license | ~13d touch | open | governance | +| 301 | RDNA1 follow-ups | ~3mo | open | stale | +| 289 | MQ3 follow-ups | ~3mo | open | partial obsolete | +| 272 | tracking/DDTree+CASK | ~3mo | open | research | +| 271 | tracking/ZAYA1 | ~3mo | open | hold | +| 270 | tracking/Gemma4 hold | ~2mo | open | supersede #672 cluster | +| 252 | bug?/Windows perf | ~3mo | open | needs repro | +| 223 | bug/string garbage | ~3mo | open | may #595 | +| 217 | bug/missing Drop | ~3mo | open | leak w/ #640 | +| 213 | bug/attractor 27B | ~3mo | open | quality | +| 209 | feat/MQ3 MoE prefill | ~3mo | open | backlog | +| 207 | feat/gfx906 MoE | ~3mo | open | backlog | +| 188 | research/mq3-lloyd | ~3mo | open | research | +| 162 | bug/NixOS install | ~3mo | open | cluster #645 | +| 155 | meta/modularization shipped | ~4mo | open | CLOSE announcement | +| 116 | research/Lloyd-MQ3 gates | ~3mo | open | stale vs V2 | +| 114 | research/MQ3 sub-9B | ~4mo | open | registry still cites | +| 113 | research/MQ3 ppl | ~4mo | open | same | +| 105 | question/CPU+GPU | ~3mo | open | answer+close | +| 92 | feat/MoE DFlash drafts | ~3mo | open | backlog | +| 89 | bug/DFlash A3B attractor | ~4mo | open | help wanted | +| 78 | feat/sliding FA | ~4mo | open | backlog | +| 77 | design/NVMe paging | ~3mo | open | design | +| 76 | design/3-tier KV | ~2mo | open | design | +| 61 | feat/gfx1151 baselines | ~1mo | open | relevant | +| 50 | test/gfx1152 | long | open | keep; rewrite CONTRIBUTING | +| 45 | roadmap/living index | ~4mo | open | refresh or close | +| 43 | roadmap/SSM persist | ~4mo | open | stale | +| 42 | roadmap/mutable hipGraph | ~3mo | open | backlog | +| 39 | research/custom DFlash train | ~4mo | open | research | +| 31 | follow-ups/PR28 | ~4mo | open | stale | +| 30 | bug/gemm_qkvza multiblock | ~3mo | open | kernels | + +Stale >30d low traffic: majority #30–#560; bulk stale label + close #155 and answered questions #433 #105 #345. + +## Confidence + +Did not re-read all docs/** (~1495). Sampled top-level owners INDEX VALIDATION BENCHMARKS MODELS ARCHITECTURE CLI GETTING_STARTED perf-checkpoints README admissions CONTRIBUTING README AGENTS CLAUDE models.json force_local daemon Cargo CI. Did not verify every issue body for already-fixed-on-master. Did not deep-diff #681 vs multi-gpu.md. PR mergeable at fetch time. Model count ~79 structural scan not jq length. JSON summary fields: broken 9 verified; missing 5; changes 8 ranked hours→week+. diff --git a/docs/plans/audit-2026-09-02/audit-Generate.md b/docs/plans/audit-2026-09-02/audit-Generate.md new file mode 100644 index 000000000..1c8d1a6c2 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-Generate.md @@ -0,0 +1,110 @@ + + +# Audit: Generate + +## Broken + +### 1. EP MoE (arch 6) generate never wired — routes to dense TP (#683) +- **path_line:** `crates/hipfire-generate/src/qwen.rs:228-252` (`5 | 6 => ep_serve_qwen35_dense_tp`) +- **verified:** true (open issue #683 + code read) +- **how known:** `generate_ep` sends arch 5 and 6 to `ep_serve_qwen35_dense_tp`, which pattern-matches only `EpArch::Qwen35DenseTp` (`qwen.rs:373-390`) and errors `EP arch mismatch (expected dense Qwen TP)`. MoE loads as `EpArch::Qwen35 { batch, .. }`. The only `EpArch::Qwen35` arm in generate is reset (`qwen.rs:740-747`). Loader admits tp≥2 MoE; generate cannot run it. +- **sibling:** `select_generation_route` EP short-circuit only names arch 9/10 (`ar.rs:882-887`); Qwen EP becomes `GenerationRoute::Unknown`. Dispatch still serves via `Unknown if m.ep.is_some()` (`ar.rs:1305-1330`). `supports_tools` excludes Unknown, so tools are refused while generation is still attempted on a broken MoE path. + +### 2. EP GPU error exits skip `ep_reset_after_abort` +- **path_line:** `qwen.rs:404-414`, `499-508`, `1014-1026`, `1193-1201`, `1464-1473`, `1579-1587` (representative) +- **verified:** true +- **how known:** Prefill/decode/download failures call `emit_active_attempt_error` then `return` with no `ep_emit_abort` / `ep_reset_after_abort`. Client-cancel paths do reset. DS4 mitigates with unconditional start-of-turn zero (`qwen.rs:882-894`); dense_tp resets at entry (`qwen.rs:320-365`). **MiniMax EP uses LCP reuse** (`qwen.rs:1410-1427`) without full per-turn reset — a mid-prefill/decode error leaves rank KV/cursors dirty for the next LCP hit (same contamination class DS4 start-of-turn already fixed). + +### 3. Vision cancel: wire terminal without GPU/DN/spec rollback +- **path_line:** `vision.rs:1007-1010`, `1380-1382`, `1418-1422` (and dots/lfm2 VL cancel sites) +- **verified:** true +- **how known:** `check_abort` / `ClientTerminalDecision::Abort` call `emit_qwen_ar_cancelled` only. Comment claims next dispatch non-zero-seq_pos reset reclaims state (`vision.rs:1005-1008`) — that is **not** attested `production_fail_closed_rollback` used by Qwen AR/DFlash/LFM (`ar.rs:3531-3532`, `dense.rs:6152-6154`). Partial VL turn leaves `seq_pos`, `conversation_tokens`, DN, KV compact, checkpoints live until a later capacity path happens to fire. + +### 4. Vision context-full DN reset is manual memset; AR uses canonical reset +- **path_line:** `vision.rs:565-617` vs `ar.rs:1948-1966` +- **verified:** true +- **how known:** AR switched to `b.dn_state.reset(gpu)` so Q8 `s_ef_residual` cannot leak (`ar.rs:1948-1952`). VL still field-wise memsets `s_matrices`/`s_scales`/`conv_states`/`s_ef_residual` and ignores reset Result — the exact incomplete pattern AR comments call unsafe. + +### 5. Bring-up dense routes: terminal Abort is empty (no GPU rollback) +- **path_line:** `dense.rs:6699-6700` (minimax), `7430-7431` (cohere), `7661-7662` (qwen2), `8363-8364` (maple); contrast LFM `6268-6270` +- **verified:** true +- **how known:** `ClientTerminalDecision::Abort => {}` suppresses success `done` but leaves GPU session state. LFM/DS4/Glimmer call `production_fail_closed_rollback`. AR bring-up comment admits the gap (`ar.rs:4531-4533`). + +### 6. `fail_closed_reset_target_and_spec` omits Maple; `reset_core_arch_key` maps 15→unknown +- **path_line:** `common.rs:1075-1182`; `ar.rs:4609-4621` +- **verified:** true +- **how known:** Reset walk covers qwen35, llama, qwen2, ds4, lfm, minimax, cohere, gemma4, glimmer — **no `m.maple_mut()`**. `reset_core_arch_key(15)` falls through to `"unknown"`. Shared fail-closed epilogue will not reset MapleState if ever reached (maple Abort currently empty so less hot, inventory still wrong). + +### 7. Grammar sample path still swallows logits download into zeros (FIX #4 sibling) +- **path_line:** `ar.rs:3432-3433`, `3909-3910`, `4079-4080`; `qwen.rs:4644-4645`, `4959-4960`, `5091-5092` +- **verified:** true +- **how known:** EP comments forbid zero-logits fallback (`qwen.rs:1091-1093`). Single-GPU AR and PP multi grammar branches still `unwrap_or_else(|_| vec![0.0; vocab_size])` → silent token-0 corruption instead of fail-closed error. + +### 8. DS4 AR/spec dispatch discards `max_think_tokens` +- **path_line:** `ar.rs:1445-1482` (`let _ = (… max_think_tokens …)` before `generate_deepseek4` / `_spec`) +- **verified:** true +- **how known:** ThinkMode is threaded; numeric think cap is not. DS4 EP decode loop also has no think-cap enforcement (parser priming only). + +### 9. Stale Maple docs claim max_think discarded; code now enforces +- **path_line:** `dense.rs:7679-7682` comment vs `dense.rs:8187-8191` + `ar.rs:1418-1441` +- **verified:** true (doc/code contradiction) +- **how known:** Module doc says dispatch discards max_think like Qwen2; MapleAr threads it and `MapleThoughtRouter` force-closes. Contributor trap. + +### 10. Qwen2 path has no cancel poll and no think/tools +- **path_line:** `dense.rs:7605-7663`; dispatch drops max_think at `ar.rs:1361-1395` +- **verified:** true +- **how known:** Decode loop never calls `check_abort`; Abort arm empty. Documented bring-up limits, still a live serve route for arch 7. + +## Missing + +### M1. No `GenerationRoute::QwenEp` / MoE EP serve +- **path_line:** `ar.rs:882-887`, `qwen.rs:240`; `tests/generation_route_matrix_tests.rs` lacks Qwen+ep row +- **verified:** true +- **summary:** Matrix covers MiniMaxEp/Deepseek4Ep only. Arch crate has `forward_ep` / batch readiness (peer AuditArchQwen); generate never calls them for single-stream MoE EP. #683 tracks wire-up; admission refusal still missing on master. + +### M2. EP reset does not clear `asst_turn_cache` / host checkpoint rings / speculator +- **path_line:** `qwen.rs:787-798` vs `common.rs:1082-1083`, `1156-1167` +- **verified:** true +- **summary:** `ep_reset_after_abort` clears seq_pos + conversation_tokens + EP GPU state only. Single-GPU production epilogue also clears asst_turn_cache and frees prefill/dflash checkpoints. EP MiniMax LCP + shared host cache can diverge after abort. + +### M3. Think-cap policy drift across routes +- **path_line:** dense_tp fail-closed `qwen.rs:561-575`; VL force-close `vision.rs:1167-1240`; AR/PP force-answer latch `qwen.rs:4777+`; Glimmer/Gemma/Maple routers force-close +- **verified:** true +- **summary:** Same user control means hard validation error (dense TP EP), force `` and continue (VL/AR/Maple), or strength-primary (Glimmer strength ignores cap at `dense.rs:2638`). No single contract table. + +### M4. Gemma4 lowered/MoE generate refused at runtime only +- **path_line:** `ar.rs:1158-1170` +- **verified:** true +- **summary:** Loader can publish lowered state; generate errors after load. Same #683-class admission gap. + +### M5. Continuous-batch eligibility reads then discards max_think +- **path_line:** `batch.rs:204-208` (`let _ = max_think`) +- **verified:** true (intentional dead read / unfinished gate) +- **summary:** Comment says 0/1/ordinary budgets are valid batch controls; gate never uses the value (budget_alert still blocks). + +### M6. No mid-decode `check_abort` on maple/qwen2/minimax/cohere bring-up loops +- **path_line:** maple/`dense.rs` decode loops; qwen2 `7605+` +- **verified:** true +- **summary:** Cancel only at commit_ready (often empty Abort). Long max_tokens burns GPU after client gone. + +## Would change + +1. **Refuse or wire Qwen MoE EP (close #683)** — `qwen.rs:240` + loader admission. Admission refuse now (**hours**); full `ep_serve_qwen35_moe` over `EpArch::Qwen35` (**days–week**). +2. **Fail-closed all EP GPU error exits through `ep_reset_after_abort` before emit** (**hours**); MiniMax post-error or start-of-turn reset when LCP unsafe (**hours**). +3. **VL cancel → `production_fail_closed_rollback` + attested cancel terminal** (match AR) (**half-day**); unify VL context-full onto `dn_state.reset` (**hours**). +4. **Extend `fail_closed_reset_target_and_spec` + `reset_core_arch_key` for maple (15)**; fill Abort arms on maple/qwen2/minimax/cohere (**day**). +5. **Replace grammar zero-logits swallow with emit+rollback** on AR/PP (copy EP FIX #4) (**hours**). +6. **Thread or explicitly refuse DS4 `max_think_tokens`** (**hours**). +7. **Add `GenerationRoute::QwenEp` (DenseTp vs Moe)** so tools/route matrix stay honest (**half-day** once serve exists). +8. **Single think-cap policy matrix test** per route (**day**). +9. **Fix stale maple module docs** (`dense.rs:7679`) (**minutes**). +10. **EP `ep_reset`: clear `asst_turn_cache`** parity with single-GPU (**hours**). + +## Confidence + +- No GPU/hw-gate/builds/tests run (read-only). +- Did not fully trace `redline.rs` capture paths or every batch lane teardown. +- Did not verify maple GPU reset API beyond missing callsite in fail_closed. +- #683 confirmed open — not re-claimed as novel; added siblings (Unknown route, tools gate, reset-only MoE arm). +- PP abort double path (`reset_pp_uncommitted_state` then `production_fail_closed_rollback`) looks intentional; not proven redundant. +- Peers own loader admission / arch EP readiness / runtime `reset_core` internals. diff --git a/docs/plans/audit-2026-09-02/audit-LoaderDaemon.md b/docs/plans/audit-2026-09-02/audit-LoaderDaemon.md new file mode 100644 index 000000000..69aadfef4 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-LoaderDaemon.md @@ -0,0 +1,5 @@ + + +# Audit: LoaderDaemon + +Client JSONL → daemon match(type) → load: tp>1 load_model_ep_with_kv_mode else load_model_with_gemma4_drafter → Carrier::load → LoadedModel{state|ep}. Post-load stage_continuous_batch may set EpArch::Qwen35.batch. Generate: set_active_attempt_id + activate_terminal_control → select_generation_route (EP only 9/10 named; 5|6 Unknown) → generate bodies. MoE EP intended serve is drive_qwen35_ep_continuous_batch only when batch staged and serve_continuous_batch. Errors: emit_active_attempt_error (TLS id) vs emit_uncorrelated_error (attempt 0) vs write_error_envelope (explicit). Unload: pflash drafter free before unload_model; ensure_vmm_ready_for_load on empty unload and pre-load. diff --git a/docs/plans/audit-2026-09-02/audit-Mq4v2Kernels.md b/docs/plans/audit-2026-09-02/audit-Mq4v2Kernels.md new file mode 100644 index 000000000..4abcb7587 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-Mq4v2Kernels.md @@ -0,0 +1,153 @@ + + +# MQ4G256V2 (qt=44) kernel-family audit — 2026-09-03 + +_Lifecycle: planned intent. Read-only audit of master `8cd15a62b`. Findings cite `path:line` at that commit._ + +**Scope.** All 34 HIP translation units that decode the qt=44 header (6 scalar decode, +21 dense WMMA prefill across gfx11/gfx12 incl. BT and MW-LDS, 6 MoE, 1 RDNA3 MMQ; ~8.4k +lines) plus the Rust encoder (`hipfire-quantize`), dispatch (`hipfire-dispatch`), +launchers (`rdna-compute`), replay contracts, and the parity examples. Five read-only +slices (scalar decode, WMMA gfx11, WMMA gfx12 + MMQ, MoE + reachability, Rust side), and +an independent pass by the auditor over every half-select site in the family. + +**Oracle.** `docs/quant-formats/mq4-v2.md` §§2–4: 136 B/group, fp16 `(scale, zero)` for +half 0 at `[0..4)` and half 1 at `[4..8)`, nibbles at `[8..136)`; the half-select +predicate must be derived from the kernel's own nibble addressing, never assumed from +lane id — a wrong predicate "compiles, runs, and silently applies the wrong scale to half +of every tensor." + +## Verdict + +**The format is sound end to end.** No wrong half-select, no scale/zero or half0/half1 +swap, no division by scale (degenerate `scale=0` reproduces `zero` exactly everywhere), +no last-group header over-read, no kernel key without a launch arm, WMMA fragment +k-order matches the dequant order on both gfx11 (full-K per lane, C rows `2j+(tid>>4)`) +and gfx12 (`k_grp` K-split, C rows `8*k_grp+j`), MMQ applies scale per 128-K half (not +a per-256 V1-ism), and Redline replay/graph capture pin the fixed base kernels rather +than the adaptive BT policy, as the spec claims. + +Half-select sites, verified against their own addressing (auditor's pass, all 34 files): + +| predicate | nibble address | files | +|---|---|---| +| `tid < 16` | `gp + 8 + tid*4` (`boff`) | 6 scalar decode, 4 MoE GEMV (`hoff` form) | +| `kt < 8` / `k_off >= 128` / `kt >= 8` | `gp + 8 + kt*8 [+ k_grp*4]` | 20 WMMA main/BT bodies, 2 MoE grouped WMMA | +| `segment < 4` | `gp + 8 + segment*16` (kt0 = 2·segment) | 3 MW-LDS stagers | +| `quarter_in_group < 2` | `gp + 8 + quarter*32 + {0,8,16,24} + k_grp*4` | 2 gfx12 ldsstage bodies | +| `ksc` slot ownership | dm slots 0..3 ← K 0..127, 4..7 ← K 128..255 | MMQ | + +MoE is **reachable and correct** on dedicated V2 kernels (loader → `MoeResolution.routed_indexable_mq4v2` +at `crates/hipfire-dispatch/src/families/moe.rs:244` → `pipeline/mod.rs:1244-1262`, +`:1460-1472`, `:1591-1607`, `:3056`, `:3291+`), with a parity harness that uses disjoint +halves and a negative control (`crates/rdna-compute/examples/mq4v2_moe_parity.rs`). + +## Broken + +1. **`hipfire_runtime::llama::is_batchable_la` is not in lockstep with `qwen35::is_batchable_la`** — high, verified. + `crates/hipfire-runtime/src/llama.rs:1878-1892` admits `MQ4G256V2` (and qt 47–50) for + WMMA prefill only on `gfx1200|gfx1201`; `crates/hipfire-arch-qwen35/src/qwen35/prefill.rs:1549-1572` + admits them on gfx11 + gfx12 via `mqv2_gfx11_wmma_enabled_from_env` (kill-switch + `HIPFIRE_MQV2_GFX11_WMMA=0`). Both doc-comments (`llama.rs:1841-1844`, `prefill.rs:1441-1442`) + claim the two "match exactly". Effect: plain Llama / Qwen3 dense models carrying qt=44 + prefill per-token on gfx1100/1151 while Qwen3.5/3.8 take WMMA. Fix is a one-line + admit plus the kill-switch, or a shared function so the lockstep is structural. + +## Missing + +1. **`mq4v2_gemm_parity` cannot discriminate a wrong half-select** — med/high, verified. + `crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs:51-63` builds Gaussian + weights (σ≈0.011) whose two halves have near-identical `(scale, zero)`, so a wrong + predicate lands inside 4-bit quantization noise — the header comment (`:19-22`) + promises a "systematic blow-up" the fixture cannot produce. `mq4v2_residual_parity.rs:15-21` + has the right construction (half 0 in `[-1,1]`, half 1 in `[96,160]`). The batch-size + sweep design in `gemm_parity` is correct and worth keeping; add a disjoint-halves arm. +2. **Production V2 kernels with no discriminating parity example:** `gemv_mq4g256v2_multirow` + (R=2/4/8, the default-R path on gfx1151/1201), `gemv_mq4g256v2_residual_sigmoid_scaled_k512` + (MoE shared-down), the MW-LDS bodies (`*_gfx1100_mw_lds`, `gemm_mqv2_wmma_gfx11_mw_lds`), + the MMQ path (`gemm_mq4g256v2_residual_mmq` — production on gfx1100/1151 at + `batch ≥ 128 && batch % 128 == 0`, `crates/rdna-compute/src/gemm.rs:17860-17909`, + call sites `:26326`, `:26815`, `:27450`, `:28035`), and the gfx1100 decode specials + (`fused_qkv_mq4g256v2_k2048_x_buffer_gfx1100`, `fused_qkvza_mq4g256v2_k2048_hoist_x32_gfx1100`). + BT is covered by `test_mq4v2_*_bt_gfx{1100,1151,1201}`. +3. **No test that `mqv2_prefill_batch_tile` (`gemm.rs:124`) only selects tiles the arch's + launcher accepts.** Policy and launchers agree today (gfx1100 BT4/12 QKV, BT6/12 gate, + BT4/6/8 residual; gfx1151 BT12 gate, BT4 else; gfx1201 BT8 QKV); drift is untested. +4. **`KernelKey::GemvMq4G256V2SwiGLUResidual` is an alias** (`families/gemv.rs:609` → + `gemv_hfq4g256_residual_mq4v2`), not a fused SwiGLU+residual kernel; the caller + pre-fuses. Registry admits a key that has no distinct symbol. +5. **No degenerate-scale (`scale bits = 0`, nonzero nibbles) fixture** through any kernel. + Arithmetic is `q*sc+zp` everywhere so it is correct by construction; a fixture would + make that a tested contract. + +## Would-change + +1. **Spec §9 (`docs/quant-formats/mq4-v2.md`) is stale in three places:** says MoE is + "out of scope / fail-closed" (it is production-wired, decode + prefill, gfx11 + gfx12); + says the XBATCH single-row path was "not ported" (`gemv_mq4g256v2.hip:295-361` has it, + correct); §4 narrative says main WMMA bodies step `kt += 4` (gfx11 bodies step by 2, + residual by 1 — still correct, each body selects with `kt < 8`). +2. **Stale in-file status comments:** `gemm_mq4g256v2_residual_mmq.hip:10` says + "Experimental" for a production fast path; `gemm_qkv_mq4g256v2_wmma.gfx12.hip:15-24,169-175` + still calls its C-map a "HYPOTHESIS"/scaffold while the sibling gate_up file + (`:30-35`) documents the same map as R9700-validated. +3. **FWHT sign generation is duplicated, and the seeds are convention:** identical LCG + (`*1103515245 + 12345`) in `hipfire-quantize` and `rdna-compute`; seeds `42`/`1042` + hardcoded at `crates/rdna-compute/src/scratch.rs:408-409` (`ensure_mq_signs`) and in + the quantizer/loader paths. Correct today; a drift in either copy is silent logit + corruption. One shared function and a named constant. +4. **Four dedicated MQ4 gfx11 BT files** (`gemm_{qkvza,qkv,gate_up,residual}_mq4g256v2_wmma_gfx11_bt.hip`) + duplicate the shared `gemm_mqv2_wmma_gfx11_bt.hip` body, which instantiates + BITS ∈ {2,3,5,6} but not 4. Likewise two MW-LDS implementations (`*_gfx1100_mw_lds` + with explicit u32 nibbles vs `gemm_mqv2_wmma_gfx11_mw_lds.hip` which also emits MQ4 + `_gfx11_mw*` symbols). A future half-select or C-map fix has to land in two places. +5. **MoE GEMV header address is lane-dependent** (`*(gp + hoff)`, `hoff = (tid<16)?0:4`) + in the four `gemv_mq4g256v2_moe_*` TUs, vs the dual-scalar-load + select form the + dense kernels use. Correct; the spec §4 notes this form turns two scalar loads into a + vector load. Perf contract, not math. +6. **BT policy keys on exact arch strings** `gfx1100|gfx1151|gfx1201`; gfx1101/1102/1150/1200 + get base WMMA with no adaptive BT. Fine if intentional; not documented as such. +7. **`gemm_mq4g256v2` "plain" GEMM is memset-Y + residual WMMA** (`gemm.rs:29030-29055`); + the name and registry entry imply a kernel that does not exist. +8. `gemv_mq4g256v2_residual.hip:92-105` macros are still named `HIPFIRE_RESIDUAL_LOAD_SC`/`_ZP` + but load the two packed half-header dwords; the v1 name invites a future "fix" that + reinterprets them as f32. `fused_qkvza_mq4g256v2.hip` carries 16+ textual copies of + the header decode across `PAIR_BUFFER`/`HYBRID_BUFFER`/default branches; all agree today. +9. gfx11/gfx12 WMMA headers decode via `(_Float16)__half2float(__ushort_as_half(...))` + (fp16→f32→fp16) where a bitcast would do; BT N-tail gathers `X` from batch 0 and + discards it instead of zero-filling. Both perf nits. + +## Confirmed by design (not findings) + +- gfx1030 with qt=44: per-token decode fallback, no WMMA (`is_batchable_la` false, + `gemm.rs:29055-29061` errors without `has_wmma()`); the spec §6 R=1 regression is not + a live path. The "gfx1030 default-R decision" in §9 can be closed as moot. +- No wave64 V2 decode path exists in any of the 34 files; every launch is wave32 + (`__launch_bounds__(32, …)`, or 32·NW). Spec §9's "wave64 unverified" is therefore + "wave64 unsupported", which is what `dtype_arch_predicate → HasWave32` enforces. +- No gfx12 MMQ for qt=44 (HFQ4 has one); gfx12 stays on fp16 WMMA. Intentional per + the RDNA3-only guard (`residual_mmq.hip:92-124` stubs otherwise). +- Encoder (`crates/hipfire-quantize/src/quant_fwht.rs:200-253`): fp16 round-trip before + quantizing, `degenerate = hi==lo || step==0 || st==0` → all `q=0`, so an fp16 + underflow of the step cannot produce `inf` codes. + +## Not read + +- The v1 (`*_hfq4g256*`) sisters were not line-diffed; the port was audited against the + spec, not against its source. +- Mixed-tier MoE TUs (`gemv_mixed_moe_*`, tags 7–18) beyond a spot-check of the + `_MQ4V2G` macro; Muse/DS4 arch crates' MoE call sites end to end. +- HIP parameter lists were not compared to the Rust kernarg blobs line by line + (representative orders recorded in the Rust slice; a mismatch would show in parity). +- Fixture contents of `test_mq4v2_*_bt_gfx*` (names and coverage only). +- Nothing was built or run; this is a source audit. Parity examples exist for every + route above except those listed under Missing 2 — running them on gfx1100/1151/1201 + is the hardware confirmation this audit does not provide. + +## Recommendation + +One code fix (Broken 1 — shared `is_batchable_la` or a matching admit), one test PR +(disjoint-halves arm in `mq4v2_gemm_parity`; parity for multirow / MW-LDS / MMQ / +sigmoid-scaled k512; a `mqv2_prefill_batch_tile` ↔ registry table test), and a docs pass +(spec §9, the two in-file status comments). None of it blocks the format; qt=44 can stay +the production default. diff --git a/docs/plans/audit-2026-09-02/audit-Quantize.md b/docs/plans/audit-2026-09-02/audit-Quantize.md new file mode 100644 index 000000000..7e920ac1c --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-Quantize.md @@ -0,0 +1,5 @@ + + +# Audit: Quantize + +{"slice":"Quantize","broken":[{"title":"MQ5G256 AWQ emit without supports_awq_sidecar","path_line":"crates/rdna-compute/src/dispatch.rs:500-532; crates/hipfire-quantize/src/pipeline.rs:5152-5171","verified":true,"summary":"Quantizer pre-scales W and emits awq_scale sidecars for MQ5G256 claiming runtime support; allowlist omits MQ5G256 so loader silently drops sidecars → (W·s)·x corruption."},{"title":"Runtime HFQ header version ignored","path_line":"crates/hipfire-runtime/src/hfq.rs:493; crates/saddle-quant/src/format/hfq.rs:82-88","verified":true,"summary":"Writer stamps version=1; runtime binds _version unused; saddle-quant rejects non-1. Future v2 would load in engine and fail in tooling."},{"title":"diag_mq4_lloyd_multiacc expects qt=21 for MQ4G256Lloyd","path_line":"crates/hipfire-runtime/examples/diag_mq4_lloyd_multiacc.rs:111-114","verified":true,"summary":"MQ4G256Lloyd renumbered 21→30; example asserts 21 and cannot open current artifacts."},{"title":"hfq_manifest qt name table wrong","path_line":"crates/hipfire-runtime/examples/hfq_manifest.rs:16-28","verified":true,"summary":"qt0 labeled F32 (is Q4F16G64); qt19 labeled MQ4G256(alt) (is MQ2G256Lloyd)."},{"title":"saddle-quant QuantType missing many register tags","path_line":"crates/saddle-quant/src/format/mod.rs:30-68","verified":true,"summary":"Missing 22,28-31,38-41 vs writer/register; inspect reports unknown; registry script does not check this enum."},{"title":"mq2 refuse message claims Lloyd not landed","path_line":"crates/hipfire-quantize/src/pipeline.rs:846-858","verified":false,"summary":"Gate enforced but error text still describes Path D Lloyd as future remediation though Lloyd formats already exist."}],"missing":[{"title":"No i-know-this-is-broken; mq2v2 ungated despite product rejection","path_line":"crates/hipfire-quantize/src/cli.rs:154-165; pipeline.rs:298; docs/quant-formats/ladder.md","verified":true,"summary":"Escape hatches are allow-mq2/lloyd/etc. mq2v2 freely writable while ladder marks mq2 product cells rejected (KLD ~12-14)."},{"title":"check-quant-registry skips saddle-quant","path_line":"scripts/check-quant-registry.py:48-53","verified":true,"summary":"Only hipfire-quantize QuantType + RAW_CODECS; cannot catch inspect-parser lag or layout/AWQ semantic drift."},{"title":"PARO qt28/29 no quantize encoder","path_line":"crates/hipfire-quantize/src/hfq.rs:139-140; qt-register.txt","verified":true,"summary":"Wire tags and arch-loaded disposition exist; hipfire-quantize has no paro format path."},{"title":"qt=46 mq4_k documented but not reserved in register","path_line":"docs/quant-formats/mq4-v2.md; docs/quant-formats/qt-register.txt","verified":true,"summary":"Hierarchical mq4_k claims qt=46 in docs without reserved row in SoT table."},{"title":"hipfire-reap research-only via quantize flags","path_line":"crates/hipfire-reap/map.md; crates/hipfire-quantize/src/cli.rs:104-115","verified":true,"summary":"No standalone CLI; --reap-overlay/--reap-bake only; status research."},{"title":"hipfire-atlas not a quant surface","path_line":"crates/hipfire-atlas/map.md; src/main.rs:1-62","verified":true,"summary":"Bench corpus JSONL tool; transitional Rust vs Python kernel_atlas.py."},{"title":"oracle format dual-binds F32 and BF16","path_line":"crates/hipfire-quantize/src/pipeline.rs:246,302","verified":false,"summary":"oracle sets both flags; F32 wins first except Cohere path — footgun."}],"changes":[{"title":"Fix MQ5G256 AWQ allowlist or stop emitting","path_line":"crates/rdna-compute/src/dispatch.rs:500","cost":"hours","summary":"Add MQ5G256 to supports_awq_sidecar or disable MQ5 AWQ emit; test encoder-qt ↔ allowlist."},{"title":"Gate mq2v2 by default","path_line":"crates/hipfire-quantize/src/pipeline.rs:298","cost":"hours","summary":"--allow-mq2v2 required; cite ladder rejection."},{"title":"Enforce HFQ version in runtime open","path_line":"crates/hipfire-runtime/src/hfq.rs:493","cost":"hours","summary":"Reject unsupported versions like saddle-quant."},{"title":"Pin saddle-quant QuantType in registry check","path_line":"scripts/check-quant-registry.py","cost":"hours-day","summary":"Third source must match quantize enum tag set."},{"title":"Fix diag + hfq_manifest qt tables","path_line":"crates/hipfire-runtime/examples/","cost":"hours","summary":"qt30 for Lloyd; names from qt-register."},{"title":"Reserve qt=46 in qt-register","path_line":"docs/quant-formats/qt-register.txt","cost":"minutes","summary":"reserved row for mq4_k hierarchical."},{"title":"Single QuantType module","path_line":"crates/saddle-quant/src/format/mod.rs","cost":"week+","summary":"One wire-tag crate; kill private enums in bins."}],"report":"(full markdown is in result.data.summary — parent persists as local://audit-Quantize.md)"} diff --git a/docs/plans/audit-2026-09-02/audit-RedlineSpec.md b/docs/plans/audit-2026-09-02/audit-RedlineSpec.md new file mode 100644 index 000000000..efa5ceb69 --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-RedlineSpec.md @@ -0,0 +1,5 @@ + + +# Audit: RedlineSpec + +Product retained replay lives in rdna-compute::replay::ReplayController (not crates/redline). Flow: daemon load → retained_redline_default (mq4r on gfx1100/1151/1201 single-GPU pp=tp=1, plus ds4 mq2r gfx1151 AR; muse_glimmer carved out) → configure_model_default arms Auto/PM4 → first eligible plain-AR forward records HIP launches → prepare_pm4/aql bakes kernarg device pointers + dynamic position/GDN bindings → state Ready → subsequent decode should_route_pm4/aql. Failures poison → sticky Fallback to HIP. Layout growth rearm_after_layout_growth clears prepared state. redline-dispatch is the hazard-checked DAG/plan library + AQL/PM4 helpers via redline-rocr; product auto path skips its shadow-install gate. BindingLayoutFingerprint deliberately omits base pointers; BindingRevision is caller-honored. radiowave supplies optional HSACO cache-class certifications and offline recipes — not live recipe selection. pflash is optional long-prefill drafter compression still linked into daemon/generate (default mode Off avoids drafter VRAM). hipfire-detect is offline JSONL observational bank (coherence harness), not in the generate hot path; thresholds partly diverge from runtime LoopGuard defaults. diff --git a/docs/plans/audit-2026-09-02/audit-Runtime.md b/docs/plans/audit-2026-09-02/audit-Runtime.md new file mode 100644 index 000000000..9063e50bb --- /dev/null +++ b/docs/plans/audit-2026-09-02/audit-Runtime.md @@ -0,0 +1,174 @@ + + +# Audit: Runtime + +# Audit — Runtime + +Slice: `crates/hipfire-runtime/src/**` (except config/loader_api/multi_gpu ownership). Focus: HFQ/qt registry vs GEMV/GEMM, free_gpu/Drop, reset_core arch keys, spec rollback, sampler/eos emit. + +## Broken + +### 1. LlamaWeights::free_gpu frees `.buf` only — AWQ/PARO sidecars leak on unload +- **path_line:** `crates/hipfire-runtime/src/llama.rs:683-706` vs `WeightTensor::free_all` at `llama.rs:525-540` and DFlash `dflash.rs:946-975` +- **verified:** true +- **how known:** Read both. `LlamaWeights::free_gpu` does `gpu.free_tensor(l.wq.buf)` (and gate/up/down/output.buf). `WeightTensor::free_all` frees `paro` pairs/theta/scales (unless alias) and `awq_scale` then `buf`. DFlash comments explicitly: "free_all (not .buf) so the awq_scale / paro sidecars are released too — .buf-only freeing leaks one tensor per weight per layer". HFQ `load_weight_tensor` attaches AWQ when `supports_awq_sidecar()` (`hfq.rs:1485-1490`). Qwen35 weights path correctly uses `free_all`. +- **impact:** Unload/reload of llama (or any consumer of `LlamaWeights::free_gpu`) with AWQ/PARO HFQ silently leaks VRAM proportional to weights×layers. + +### 2. weight_gemv builds incomplete WeightRef (row_stride=0, awq_scale=None, rotation=None) +- **path_line:** `crates/hipfire-runtime/src/llama.rs:740-748` vs correct adapter `llama.rs:545-560` (`dispatch_ref`) +- **verified:** true +- **how known:** Read `weight_gemv`: hardcodes `row_stride: 0`, `rotation: None`, `awq_scale: None` while `WeightTensor::dispatch_ref` wires `row_stride`, Givens, and `awq_scale`. Paro path re-reads `w.paro` separately; MQ rotation uses `rotate_x_mq_for(gpu, w, …)` so AWQ-on-activation may still apply there — but any kernel path that reads AWQ/row_stride from `WeightRef` (MQ4C padded stride, weight-side AWQ) gets wrong metadata. GL/MQ4C/dense AWQ consumers of this entry point are at risk. +- **impact:** Silent wrong decode or dropped sidecars on formats that need non-zero `row_stride` or weight-ref AWQ; dead `dispatch_ref` helper. + +### 3. MQ2G256GL / MQ3G256GL load as passthrough but have no dense GEMV/prerotated key +- **path_line:** `weight_backend.rs:411-422` (RAW_CODECS qt 38/39); `hipfire-dispatch/src/types.rs:120-125` (FwhtG256); `types.rs:730-776` (`for_gemv_prerotated` Err for rotation dtypes not listed); tests `hipfire-dispatch/src/tests.rs:1348-1365` (`gl_dtypes_have_no_dense_gemv_key`) +- **verified:** true +- **how known:** RAW_CODECS admits 38/39. Tests assert plain and prerotated GEMV keys are Err (MoE-indexed only). `weight_gemv` residual arm for FwhtG256 rotates then `GemvVariant::Prerotated` → `for_gemv_prerotated` → UnsupportedVariant. Loader admits tensors the dense engine cannot run (archetype of #683). +- **impact:** Dense GL weights fail at first GEMV; or if any fallthrough mis-routes, garbage. MoE-indexed path is the only supported use. + +### 4. HFQ `load_weight_tensor` vs `dequant_weight_raw` F16/F32/BF16 divergence +- **path_line:** `hfq.rs:1456-1483` (qt1→host F32 only; else raw_codec or error); `weight_backend.rs:681-724` (qt1 keep F16; qt2 F32; qt16 BF16→F32) +- **verified:** true +- **how known:** Comments in both files document divergence. HFQ path never special-cases qt 2 or 16 → `raw_codec` None → "unsupported quant_type". Safetensors/backend path keeps F16 on device for qt1 (native GemvF16) while HFQ forces F32 GEMV. Same on-disk qt, different GPU dtype and kernel family. +- **impact:** F32/BF16 weight tensors in HFQ fail load; F16 HFQ models pay 2× weight bandwidth and lose F16 GEMM/GEMV path vs ST path. + +### 5. dots_ocr arch_key spelling contradicts reset_core inventory and generate map +- **path_line:** `hipfire-arch-dots-ocr/src/arch_model.rs:24-26` (`"dots_ocr"`); `reset_core.rs:140-152` (`arch: "dots-ocr"`); `hipfire-generate/src/ar.rs:4608-4628` (`8 => "dots-ocr"`); contract comment `arch_model.rs:70-72` ("Matches the key used by reset_core … so the two cannot drift") +- **verified:** true +- **how known:** Cross-read three sites. `reset_coverage_for(bundle.arch_key())` → None for dots_ocr; `reset_coverage_for("dots-ocr")` → Some(Ineligible). Today both yield non-eligible, so retry behavior is coincidentally same; coverage checklist and any future Eligible flip or graph flags are invisible via ArchModel path. Contract comment is violated. + +### 6. generate `reset_core_arch_key` drops maple (arch_id 15 → "unknown") +- **path_line:** `hipfire-generate/src/ar.rs:4608-4628`; maple `arch_key` `hipfire-arch-maple/.../bundle.rs:40-42` (`"maple"`); inventory `reset_core.rs:196-220` (MAPLE row present) +- **verified:** true +- **how known:** match arms stop at 14 muse_glimmer; `_ => "unknown"`. Inventory has maple; ArchModel returns maple; arch_id path never sees the row. + +### 7. DFlash target_hidden row gap: warn once, still poison session +- **path_line:** `dflash.rs:1083-1100` (GAP_WARNED atomic; eprintln; continues); comments 1083-1088 (NaN hole collapses acceptance for rest of session, survives prompt-cache HIT) +- **verified:** true (code path exists; runtime trigger depends on caller advancing without hidden commit) +- **how known:** Read TargetHiddenLog commit path. Documented class (#462 / forced-advance holes). Mitigation is warn-not-fail; `spec.rs:743-747` documents forced-advance must fill hidden or leave hole. +- **impact:** Silent acceptance collapse until reseed; easy to miss in logs (warn-once). + +### 8. Generic DFlash Speculator ignores SpecRequestConfig fields (min_p / ngram / full seed contract) +- **path_line:** `dflash_generic.rs:~1000-1015` (comments: New SpecRequestConfig fields min_p/rng_seed/ngram ignored; only temp/top_p/top_k/rng_state applied) +- **verified:** true +- **how known:** Read configure_request. Documented incomplete vs qwen35 DflashSpeculator. +- **impact:** LlamaCarrier generic DFlash sampling/ngram behavior diverges from request config; user thinks knobs work. + +## Missing + +### 1. No load-time refusal for passthrough qts without dense GEMV (GL and similar) +- **path_line:** RAW_CODECS + absence of gate in `decode_raw_codec` / model_load +- **verified:** true (gap) +- Loader admits MoE-only formats into WeightTensor used by dense `weight_gemv`/`weight_gemm` without tagging or checking consumer. Missing: disposition `moe-indexed-only` in qt-register + load assert or dtype capability bit. + +### 2. MFP2/MFP3 E8 (and other arch-loaded Fwht dtypes) missing from `for_gemv_prerotated` explicit arms +- **path_line:** `hipfire-dispatch/src/types.rs:730-776` (lists MFP4* E8 but not MFP3G32E8/MFP2G32E8); rotation plan marks them FwhtG256 (`types.rs:124-125`) +- **verified:** true for key table; arch crates may use private gemv_auto +- Dense `weight_gemv` prerotated path would Err. Arch-loaded only in register (`qt-register.txt:36-37`) — OK if never routed through llama weight_gemv; missing shared capability matrix so contributors do not wire them into dense path by accident. + +### 3. weight_gemm only special-cases a subset; rest is per-row GEMV fallback (incl. formats that then fail) +- **path_line:** `llama.rs:1475-1575` (`_ =>` loop weight_gemv) +- **verified:** true +- Missing batched GEMM for TQ2/BQ1/Lloyd/MFP/GL/etc. Fallback multiplies cost and inherits weight_gemv gaps (GL hard-fail). + +### 4. reset_core inventory vs ArchModel arch_key not mechanically enforced for all arches +- **path_line:** `reset_core.rs` tests pin inventory strings; `arch_model.rs` comments require match; dots_ocr proves drift landed +- **verified:** true (test gap) +- Missing compile-time or CI check that every `ArchModel::arch_key` equals inventory / `reset_core_arch_key(arch_id)`. + +### 5. EosFilter is solid post-stop; missing unified guarantee that all generate loops never emit after terminal without filter +- **path_line:** `eos_filter.rs:34-49`, `186-189` (stopped → Hold); module docs say daemon loops decode+ship +- **verified:** filter correct; full daemon wire audit not completed in this slice +- EmitAndStop contract (marker not in payload) is implemented. Residual risk is loops that bypass EosFilter or emit raw detok after Stop — hand off to Generate/Cli scouts. + +### 6. Maple/Qwen MTP free_gpu still `.buf` on some WeightTensors (cross-slice) +- **path_line:** e.g. arch-qwen35 `mtp_head.rs` frees `eh_proj.buf` / expert `.buf` (seen in grep); Qwen35 main weights use free_all +- **verified:** true for MTP head pattern; primary LlamaWeights issue is in-slice +- Half-migration: free_all adopted in dflash/qwen35 trunk, not universal. + +## Would change (ranked) + +1. **LlamaWeights::free_gpu → free_all for all WeightTensors** (and audit every `free_tensor(.*\.buf)` on WeightTensor owners in runtime) + - path: `llama.rs:683-706` + - cost: **hours** + - Fix leak class DFlash already documented; add unload smoke that counts device allocs with AWQ model. + +2. **weight_gemv/residual/swiglu: use `w.dispatch_ref()` only** + - path: `llama.rs:740+` + - cost: **hours** + - Delete hand-rolled WeightRef; single wire for row_stride/AWQ/Givens. Unit test MQ4C row_stride and AWQ flag propagation. + +3. **Refuse or tag MoE-only dtypes at load (GL 38/39)** + - path: `weight_backend.rs` + `qt-register.txt` + optional `DType::dense_gemv_supported` + - cost: **hours–1 day** + - Fail fast with clear error instead of first-token UnsupportedVariant; align register disposition. + +4. **Unify HFQ vs dequant host-decode for qt 1/2/16** + - path: `hfq.rs:1456+`, `weight_backend.rs:681+` + - cost: **1 day** + - Prefer keep-F16 for qt1 (match dequant + GemvF16); add qt2/16 arms to HFQ loader or route HFQ through dequant_weight_raw exclusively. + +5. **Normalize arch_key spellings (dots_ocr ↔ dots-ocr) + extend reset_core_arch_key for maple** + - path: dots-ocr arch_model, reset_core, generate ar.rs + - cost: **hours** + - Pick one SoT string; CI test: ∀ ArchModel key ∈ inventory keys and ∀ arch_id map value ∈ inventory. + +6. **TargetHiddenLog gap: hard error or auto-reseed, not warn-once** + - path: `dflash.rs:1083-1100` + - cost: **1 day** (call-site fixes for forced-advance may be more) + - Prevent silent session poison; pair with Generate scout on forced-suffix paths. + +7. **Generic DFlash configure_request parity with SpecRequestConfig** + - path: `dflash_generic.rs` + - cost: **hours–1 day** + - Wire min_p/ngram/seed or document unsupported and refuse non-default. + +8. **Capability matrix: QuantType × load path × dense GEMV × GEMM × MoE × embed** + - path: docs/quant-formats + scripts/check-quant-registry.py extension + - cost: **days** + - Prevent #683-class (loader admits, engine cannot run) systematically; include GL, LloydU, TQ2/BQ1, arch-loaded E8. + +## Confidence + +**Did:** Read qt-register, RAW_CODECS full table, dequant_weight_raw, hfq load_weight_tensor, LlamaWeights free_gpu vs free_all vs dflash free_all, weight_gemv/weight_gemm, for_gemv_prerotated + GL tests, reset_core inventory + generate arch_id map, dots_ocr/maple arch_key, TargetHiddenLog gap, EosFilter Stop/Hold, spec commit_prefix contracts, generic dflash configure_request. + +**Did not fully:** Every residual/swiglu WeightRef twin; full dspark_core position rewind line-by-line; every daemon emit loop vs EosFilter (Generate/Cli); KvCache VMM free edge cases beyond free_gpu tests; safetensors_source upload vs mmap pager in depth; open GitHub issue cross-check via `gh` (no shell in this scout); GPU/runtime reproduction. + +**Verified vs suspicious:** Items 1–6 and 8 verified by code read/cross-ref. Item 7 (hidden gap) verified as warn-continue behavior; actual production triggers need caller trace. MFP2/3 prerotated gap is real in dispatch table; may be intentional if only arch private paths use them — treat as missing guardrails unless a dense callsite is found. + +## JSON summary (for parent merge) + +```json +{ + "slice": "Runtime", + "broken": [ + {"title": "LlamaWeights free_gpu leaks AWQ/PARO sidecars", "path_line": "crates/hipfire-runtime/src/llama.rs:683-706", "verified": true, "summary": "Frees .buf only; free_all exists and dflash uses it; AWQ attach on HFQ load."}, + {"title": "weight_gemv incomplete WeightRef", "path_line": "crates/hipfire-runtime/src/llama.rs:740-748", "verified": true, "summary": "row_stride=0, awq_scale=None, rotation=None; dispatch_ref unused."}, + {"title": "GL qts load without dense GEMV", "path_line": "crates/hipfire-runtime/src/weight_backend.rs:411-422", "verified": true, "summary": "qt 38/39 RAW_CODECS; for_gemv_prerotated Err; MoE-only by design."}, + {"title": "HFQ vs dequant F16/F32/BF16 split", "path_line": "crates/hipfire-runtime/src/hfq.rs:1456-1483", "verified": true, "summary": "HFQ qt1→F32, rejects 2/16; dequant keeps F16 and handles 2/16."}, + {"title": "dots_ocr arch_key drift", "path_line": "crates/hipfire-arch-dots-ocr/src/arch_model.rs:24-26", "verified": true, "summary": "ArchModel dots_ocr vs inventory/generate dots-ocr."}, + {"title": "maple arch_id map missing", "path_line": "crates/hipfire-generate/src/ar.rs:4608-4628", "verified": true, "summary": "arch_id 15 → unknown; inventory has maple."}, + {"title": "DFlash target_hidden gap warn-not-fail", "path_line": "crates/hipfire-runtime/src/dflash.rs:1083-1100", "verified": true, "summary": "Documents NaN poison; warns once and continues."}, + {"title": "Generic DFlash ignores SpecRequestConfig fields", "path_line": "crates/hipfire-runtime/src/dflash_generic.rs:1000-1015", "verified": true, "summary": "min_p/ngram ignored; temp/top_p/top_k only."} + ], + "missing": [ + {"title": "Load-time refusal for MoE-only passthrough qts", "path_line": "crates/hipfire-runtime/src/weight_backend.rs:317-470", "verified": true, "summary": "No dense_gemv capability check at decode_raw_codec."}, + {"title": "MFP2/3 not in for_gemv_prerotated", "path_line": "crates/hipfire-dispatch/src/types.rs:730-776", "verified": true, "summary": "FwhtG256 but no prerotated arm; arch-loaded only."}, + {"title": "weight_gemm incomplete batched coverage", "path_line": "crates/hipfire-runtime/src/llama.rs:1475-1575", "verified": true, "summary": "Many dtypes fall back to per-row GEMV."}, + {"title": "CI arch_key ↔ inventory lock", "path_line": "crates/hipfire-runtime/src/reset_core.rs:229-244", "verified": true, "summary": "dots_ocr drift proves comment-only contract."}, + {"title": "Cross-loop post-terminal emit audit", "path_line": "crates/hipfire-runtime/src/eos_filter.rs:186-189", "verified": false, "summary": "Filter OK; daemon bypass paths not fully traced here."}, + {"title": "Universal free_all on WeightTensor owners", "path_line": "crates/hipfire-runtime/src/llama.rs:525-540", "verified": true, "summary": "free_all exists; not all free_gpu call sites use it."} + ], + "changes": [ + {"title": "LlamaWeights free_all cutover", "path_line": "crates/hipfire-runtime/src/llama.rs:683", "cost": "hours", "summary": "Match dflash; AWQ unload smoke."}, + {"title": "weight_gemv use dispatch_ref", "path_line": "crates/hipfire-runtime/src/llama.rs:740", "cost": "hours", "summary": "One WeightRef construction path."}, + {"title": "Refuse GL as dense weights", "path_line": "crates/hipfire-runtime/src/weight_backend.rs:411", "cost": "hours", "summary": "Fail at load or tag moe-only."}, + {"title": "Unify qt1/2/16 load paths", "path_line": "crates/hipfire-runtime/src/hfq.rs:1456", "cost": "1 day", "summary": "Keep F16; accept F32/BF16 on HFQ."}, + {"title": "arch_key spelling + maple map", "path_line": "crates/hipfire-runtime/src/reset_core.rs:140", "cost": "hours", "summary": "Single string SoT + CI."}, + {"title": "Hard-fail target_hidden gaps", "path_line": "crates/hipfire-runtime/src/dflash.rs:1083", "cost": "1 day", "summary": "Stop silent poison."}, + {"title": "Generic DFlash SpecRequestConfig parity", "path_line": "crates/hipfire-runtime/src/dflash_generic.rs:1000", "cost": "hours", "summary": "Wire or refuse."}, + {"title": "Quant capability matrix in registry check", "path_line": "docs/quant-formats/qt-register.txt", "cost": "days", "summary": "Loader×engine matrix gate."} + ], + "report": "inline-in-architecture-field" +} +```