diff --git a/AGENTS.md b/AGENTS.md index 3e0288765..163deed66 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -220,27 +220,27 @@ works, what to measure, what counts as pass/fail. ### Pull the model + draft you want to test -Targets and drafts are independent pulls — drafts auto-discover their -target by filename when the daemon loads: +`hipfire pull ` fetches the target plus its registry-declared +DFlash draft sidecar (same mechanism as the MTP/DSpark sidecars): ```bash # 27B Qwen 3.5 (the canonical perf-test target): -hipfire pull qwen3.5:27b # 15 GB target -hipfire pull qwen3.5:27b-draft # 0.92 GB DFlash draft +hipfire pull qwen3.5:27b # 15 GB target + 0.92 GB DFlash draft sidecar # 27B Qwen 3.6 (refresh): -hipfire pull qwen3.6:27b # 15 GB target -hipfire pull qwen3.6:27b-draft # 0.92 GB DFlash draft +hipfire pull qwen3.6:27b # 15 GB target + 0.92 GB DFlash draft sidecar # 9B Qwen 3.5 (smaller, faster sanity-check): -hipfire pull qwen3.5:9b # 5.3 GB target -hipfire pull qwen3.5:9b-draft # 0.55 GB DFlash draft +hipfire pull qwen3.5:9b # 5.3 GB target + 0.55 GB DFlash draft sidecar ``` -Files land at `~/.hipfire/models/` matching the -daemon's auto-discovery pattern (`qwen3{ver}-{size}-dflash-{quant}.hfq`). -**Do not rename.** Renaming breaks the auto-discovery and DFlash falls -back to AR silently. +Standalone `*-draft` tags (`hipfire pull qwen3.5:27b-draft`) still work — +they address the same file for anyone who wants the draft alone. + +Files land at `~/.hipfire/models/`. +**Do not rename.** Load resolves the draft by its registry-declared +filename; renaming breaks the pairing — `dflash_mode auto` then runs AR +(one warning line), `on` fails the load. ### Verify md5s after pull (paranoid mode) @@ -330,7 +330,8 @@ Standalone: `cargo run --release -p hipfire-runtime --example encode_prompt -- M - `hipfire-models/qwen3.6-27b/qwen36-27b-dflash-mq4.hfq` (+ the 3.6 27B target `hipfire-models/qwen3.6-27b/qwen3.6-27b.mq4`) -Pullable via `hipfire pull qwen3.{5,6}:{9b,27b}-draft` and `hipfire pull qwen3.6:27b`. +`hipfire pull ` fetches the target plus its draft sidecar; +standalone drafts stay pullable via `hipfire pull qwen3.{5,6}:{9b,27b}-draft`. --- @@ -355,6 +356,7 @@ hipfire bench --runs 5 --warmups 3 --max-tokens 128 --json | `--spec` | `off`/`dflash`/`mtp`/`ngram`/`dspark`/`auto` | | `--backend` | `noslots` (sequential daemon) / `slots` / `batch` / `both` | | `--workload` | `stateless` / `multiturn` / `both` | +| `--prompt-file PATH` | verbatim prompt bytes for the run; JSON records `prompt_tokens`/`prompt_md5`/`prompt_chars` plus a `warnings` caveat below 256 tokens | | `--kv-mode`, `--kv-backend` | KV format and allocator | | `--reasoning-on` | off by default: a reasoning model cannot close `` inside the token budget, and the daemon fails that turn closed | @@ -429,17 +431,17 @@ reassurance. If you're testing an actual user UX flow: ```bash -hipfire pull qwen3.5:9b -hipfire pull qwen3.5:9b-draft +hipfire pull qwen3.5:9b # target + draft sidecar in one pull hipfire config set dflash_mode auto # opt in (default since 2026-04-26: off) hipfire run qwen3.5:9b "Write a Python function to find the longest substring without repeating characters" -# expected: daemon logs '[hipfire] DFlash draft detected: ...' +# expected: loader logs 'DFlash draft loaded: ...' # response generates at ≥250 tok/s on a 9B target with a paired draft ``` Without the `dflash_mode auto` config, `hipfire run` runs pure AR -even when a paired draft is on disk — the daemon explicitly logs -`[hipfire] DFlash disabled (dflash_mode=off).` This is the "I pulled +even when a paired draft is on disk. `dflash_mode on` instead requires +the sidecar and fails the load when it is missing; `developer.dflash_draft` +or `run --model-draft` overrides the sidecar. This is the "I pulled the draft but DFlash isn't firing" pitfall. --- @@ -584,8 +586,8 @@ against the A3B MoE DFlash perfmaxx line. | 3.6-A3B DFlash 68.6 tok/s vs AR 135 tok/s (50% loss) | 3.6 draft trained on 3.5 traces; target distribution mismatch on code. τ=1.22 on hard code. | Use AR mode for 3.6-A3B. Draft mismatch is expected and no 3.6 retrain is planned — Path C (`feat/mtp-dflash-training`) is dead/out-of-scope, not a forthcoming fix. 3.5-A3B DFlash works (τ=4.91). | | `hipMalloc out of memory` at hidden_rb | Long ctx (≥16K real tokens) + 27B + asym3 = tight on 24 GB | Reduce ctx, use a smaller target, or wait for the bounded-rolling-buffer trick (roadmap) | | `tok/s` below expected on long-ctx | KV cache growth — prefill is fine but decode slows past ~2K | Test at small ctx first, then scale | -| daemon doesn't auto-find draft | Filename doesn't match `qwen3{ver}-{size}-dflash-{quant}.hfq` | Don't rename the file after pull | -| `[hipfire] DFlash disabled (dflash_mode=off)` | Default flipped to `off` in 35265c6 (post-2026-04-26). Pulling a draft does NOT auto-enable DFlash anymore. | `hipfire config set dflash_mode auto` (or `on`); or per-model `hipfire config qwen3.5:9b set dflash_mode on` | +| daemon doesn't pair a pulled draft | Renamed draft file, or pulled before the sidecar existed | Don't rename files after pull; re-run `hipfire pull ` to fetch the registry-declared sidecar | +| `[hipfire-daemon] dflash_mode=off — skipping draft load` | Default flipped to `off` in 35265c6 (post-2026-04-26). Pulling a draft does NOT auto-enable DFlash anymore. | `hipfire config set dflash_mode auto` (or `on`); or per-model `hipfire config qwen3.5:9b set dflash_mode on` | | "Numbers don't match the README" | Forgot `HIPFIRE_NORMALIZE_PROMPT=1` (pre-2026-04-26) | Now default ON. Pull latest. If you opted out via `prompt_normalize=false`, that overrides the default — flip back. | | "27B DFlash regressed 30-40% suddenly" | PR #32 (cleanup-dead-wmma-kernels) on master removed `gemm_hfq4g256_residual_wmma{,2,_k4}.hip` thinking dead. Dispatch fell back to slower variants. | Verify against canonical 199 tok/s @ max=120 with default flags. If kernel files missing in `kernels/src/`, `git checkout` from a known-good commit (see commit 9a2c667 for the full recovery context). | | `HIPFIRE_GRAPH=1` reports plausible tok/s but output is garbage | Dangling stack-pointer kernargs from raw `self.hip.launch_kernel(...)` calls in `forward_scratch_layers` (kv_cache_write_*, attention_flash_*, fused_qkv_hfq4g256, rmsnorm_batched, rope_partial_interleaved_f32, gated_delta_net_q8, etc.) — captured pointers dangle past `end_graph_capture` | Bench tok/s alone never proves graph correctness. Always coherence-gate or eyeball under `HIPFIRE_GRAPH=1`. Fix: migrate every raw-launch helper used in forward_scratch_layers to `launch_maybe_blob` (model after `conv1d_silu_split_f32_n`). | @@ -602,7 +604,7 @@ against the A3B MoE DFlash perfmaxx line. | `HIPFIRE_PROMPT_HEAT_LIMIT` | Max rows in heat dump | 64 | | `HIPFIRE_KV_MODE` | Override kv_cache config | (config) | | `HIPFIRE_ATTN_FLASH` | Override flash_mode config | (config) | -|`HIPFIRE_DFLASH_DRAFT`|Force a specific draft path. Empty string = explicit opt-out|(filename auto-match alongside target)| +|`HIPFIRE_DFLASH_DRAFT`|Force a specific draft path, overriding the registry sidecar. Empty string = explicit opt-out|(unset: registry sidecar when `dflash_mode` is `auto`/`on`)| |`HIPFIRE_DFLASH_CTX_CAP`|Max rows for draft context-indexed structures (target_hidden, draft K/V caches, hidden ring). Bounds draft-side VRAM on large-`max_seq` serve loads; over-cap requests fall back to AR (identical output, slower). `0` = uncapped legacy.|8192| |`HIPFIRE_DFLASH_WINDOW`|Windowed draft context (NInfer pattern): SWA over the last W rows on draft layers 0..n-2 + full-attention last layer reaching min(physical_cap, 4W). Draft VRAM pins at W regardless of `max_seq`; past-W requests degrade τ instead of falling back to AR. Refused with CASK eviction. `0`/unset = Legacy (cap + AR fallback).|0 (off)| | `HIPFIRE_LM_HEAD_F16` | `auto`/`native` keeps qt=1 lm_head as F16; `f32`/`legacy` expands to F32 | auto/native | diff --git a/Cargo.lock b/Cargo.lock index 4ed89d3d8..53e88308f 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1313,6 +1313,7 @@ dependencies = [ "hyper", "hyper-util", "libc", + "md5", "saddle-core", "serde", "serde_json", diff --git a/crates/hipfire-arch-gemma4/map.md b/crates/hipfire-arch-gemma4/map.md index 99e65166a..8f9875990 100644 --- a/crates/hipfire-arch-gemma4/map.md +++ b/crates/hipfire-arch-gemma4/map.md @@ -24,26 +24,26 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| | [`src/arch.rs`](src/arch.rs) | 77 | 2 | 1 | -| [`src/carrier.rs`](src/carrier.rs) | 203 | 4 | 1 | +| [`src/carrier.rs`](src/carrier.rs) | 261 | 6 | 2 | | [`src/config.rs`](src/config.rs) | 880 | 19 | 12 | | [`src/drafter.rs`](src/drafter.rs) | 1,074 | 15 | 0 | | [`src/forward.rs`](src/forward.rs) | 2,665 | 6 | 4 | | [`src/gemma4.rs`](src/gemma4.rs) | 1,088 | 13 | 0 | | [`src/gemma4_vision.rs`](src/gemma4_vision.rs) | 16 | 3 | 0 | -| [`src/lib.rs`](src/lib.rs) | 48 | 8 | 0 | +| [`src/lib.rs`](src/lib.rs) | 51 | 8 | 0 | | [`src/lowered.rs`](src/lowered.rs) | 5,876 | 35 | 0 | | [`src/speculative.rs`](src/speculative.rs) | 252 | 6 | 0 | ### Public API surface - [`src/arch.rs`](src/arch.rs): `ARCH_ID`, `Gemma4` -- [`src/carrier.rs`](src/carrier.rs): `Gemma4EagerBundle`, `Gemma4LoweredBundle`, `Gemma4Bundle`, `load_gemma4_bundle` +- [`src/carrier.rs`](src/carrier.rs): `LOWERED_GENERATE_REFUSAL`, `gemma4_lowered_refusal`, `Gemma4EagerBundle`, `Gemma4LoweredBundle`, `Gemma4Bundle`, `load_gemma4_bundle` - [`src/config.rs`](src/config.rs): `LayerType`, `RopeType`, `Gemma4ESeriesVariant`, `Gemma4Config`, `from_hfq`, `from_metadata_json`, `n_full_layers`, `n_sliding_layers`, `n_full_kv_slots`, `n_sliding_kv_slots`, `max_head_dim`, `max_q_dim`, +7 more - [`src/drafter.rs`](src/drafter.rs): `DRAFTER_ARCH_ID`, `Gemma4DrafterConfig`, `from_hfq`, `pre_proj_in`, `max_q_dim`, `max_head_dim`, `DrafterLayerWeights`, `Gemma4DrafterWeights`, `load`, `free_gpu`, `Gemma4DrafterScratch`, `new`, +3 more - [`src/forward.rs`](src/forward.rs): `decode_step`, `decode_step_capture`, `decode_step_with_graph`, `supports_batched_prefill`, `forward_batch`, `forward_batch_spec` - [`src/gemma4.rs`](src/gemma4.rs): `SlidingLayerWeights`, `FullLayerWeights`, `LayerWeights`, `PerLayerBranchWeights`, `PerLayerInputWeights`, `Gemma4Weights`, `load`, `free_gpu`, `Gemma4State`, `new`, `new_with_max_seq`, `new_with_fwht3_max_seq`, +1 more - [`src/gemma4_vision.rs`](src/gemma4_vision.rs): `Gemma4VisionConfig`, `Gemma4VisionWeights`, `Gemma4VisionScratch` -- [`src/lib.rs`](src/lib.rs): `arch`, `config`, `drafter`, `forward`, `gemma4`, `lowered`, `speculative`, `carrier` +- [`src/lib.rs`](src/lib.rs): `arch`, `carrier`, `config`, `drafter`, `forward`, `gemma4`, `lowered`, `speculative` - [`src/lowered.rs`](src/lowered.rs): `wmma_prefill_enabled`, `batched_prefill_enabled`, `LayerType`, `RopeType`, `Gemma4Config`, `config_from_hfq`, `SlidingLayerWeights`, `FullLayerWeights`, `MoeExpertWeights`, `MoeLayerExtras`, `LayerWeights`, `Gemma4Weights`, +23 more - [`src/speculative.rs`](src/speculative.rs): `SpecStepOut`, `Gemma4SpecScratch`, `new`, `set_seed_hidden_from`, `free`, `spec_step_gemma4_eagle` @@ -60,6 +60,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 10 modules · 12,179 lines · 111 public items · 18 tests · 8 examples +- 10 modules · 12,240 lines · 113 public items · 19 tests · 8 examples diff --git a/crates/hipfire-arch-gemma4/src/carrier.rs b/crates/hipfire-arch-gemma4/src/carrier.rs index 4cbbc91af..ef503e91f 100644 --- a/crates/hipfire-arch-gemma4/src/carrier.rs +++ b/crates/hipfire-arch-gemma4/src/carrier.rs @@ -39,6 +39,29 @@ fn gemma4_validate_drafter_route(is_e_series: bool, has_drafter: bool) -> Result Ok(()) } +/// Generate-time refusal for the lowered/MoE path, shared with the load-time +/// admission below so both name the same combination and remedy. The generate +/// body is eager-dense-only; a lowered load must never reach it. +pub const LOWERED_GENERATE_REFUSAL: &str = "gemma4 lowered/MoE generate not yet wired on this build (eager dense only) — reload without batched/WMMA prefill opt-in or the MoE variant"; + +/// Admission decision for the lowered path, before any device allocation. +/// Pure so the contract is unit-testable. Returns the refusal reason when the +/// (model, option) combination would select `lowered` via [`gemma4_use_lowered`], +/// which generate cannot serve — fail here instead of after a full weight/KV +/// upload. `None` means the eager dense path serves the combination. +pub fn gemma4_lowered_refusal( + enable_moe_block: bool, + want_batched: bool, + has_drafter: bool, + is_e_series: bool, +) -> Option<&'static str> { + if gemma4_use_lowered(enable_moe_block, want_batched, has_drafter, is_e_series) { + Some(LOWERED_GENERATE_REFUSAL) + } else { + None + } +} + // ─── Bundle types ───────────────────────────────────────────────────────── pub struct Gemma4EagerBundle { @@ -119,6 +142,22 @@ pub fn load_gemma4_bundle(src: ModelSource, ctx: &mut LoadCtx) -> Result>(); assert_eq!(lowered_kv_layer_counts(&layer_types), (40, 8)); } + + #[test] + fn lowered_admission_refuses_what_generate_cannot_serve() { + // MoE variant: refused with the combination and remedy named. + let err = gemma4_lowered_refusal(true, false, false, false) + .expect("MoE lowered must be refused at admission"); + assert!(err.contains("lowered/MoE"), "reason names the state: {err}"); + assert!(err.contains("eager dense only"), "reason: {err}"); + // Dense with batched/WMMA prefill opt-in and no drafter: refused. + assert!(gemma4_lowered_refusal(false, true, false, false).is_some()); + // Adjacent supported: eager dense (no opt-in), drafter-kept-eager, + // and E-series all admit. + assert_eq!(gemma4_lowered_refusal(false, false, false, false), None); + assert_eq!(gemma4_lowered_refusal(false, true, true, false), None); + assert_eq!(gemma4_lowered_refusal(false, true, false, true), None); + // MoE always selects lowered even with a drafter requested, so it is + // refused too (the drafter route only keeps *dense* eager). + assert!(gemma4_lowered_refusal(true, false, true, false).is_some()); + } } diff --git a/crates/hipfire-arch-gemma4/src/lib.rs b/crates/hipfire-arch-gemma4/src/lib.rs index 23d2e83f0..cd8dfaecc 100644 --- a/crates/hipfire-arch-gemma4/src/lib.rs +++ b/crates/hipfire-arch-gemma4/src/lib.rs @@ -28,14 +28,17 @@ //! `gelu_tanh_f32`, `logit_softcap_f32`, plus the shared GEMV path. pub mod arch; +pub mod carrier; pub mod config; pub mod drafter; pub mod forward; pub mod gemma4; pub mod lowered; pub mod speculative; -pub mod carrier; -pub use carrier::{load_gemma4_bundle, Gemma4Bundle, Gemma4EagerBundle, Gemma4LoweredBundle}; +pub use carrier::{ + gemma4_lowered_refusal, load_gemma4_bundle, Gemma4Bundle, Gemma4EagerBundle, + Gemma4LoweredBundle, LOWERED_GENERATE_REFUSAL, +}; pub use arch::{Gemma4, ARCH_ID}; pub use config::{Gemma4Config, LayerType, RopeType}; diff --git a/crates/hipfire-arch-qwen2/map.md b/crates/hipfire-arch-qwen2/map.md index 87ee2eb77..1fc4c7d8a 100644 --- a/crates/hipfire-arch-qwen2/map.md +++ b/crates/hipfire-arch-qwen2/map.md @@ -26,7 +26,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | [`src/arch_model.rs`](src/arch_model.rs) | 45 | 0 | 0 | | [`src/carrier.rs`](src/carrier.rs) | 76 | 2 | 0 | | [`src/lib.rs`](src/lib.rs) | 85 | 5 | 0 | -| [`src/qwen2.rs`](src/qwen2.rs) | 2,355 | 23 | 7 | +| [`src/qwen2.rs`](src/qwen2.rs) | 2,358 | 23 | 7 | | [`src/spec_impl.rs`](src/spec_impl.rs) | 199 | 1 | 0 | ### Public API surface @@ -51,6 +51,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 6 modules · 2,849 lines · 32 public items · 8 tests · 4 examples +- 6 modules · 2,852 lines · 32 public items · 8 tests · 4 examples diff --git a/crates/hipfire-arch-qwen2/src/qwen2.rs b/crates/hipfire-arch-qwen2/src/qwen2.rs index 0535c2556..7ab2054d7 100644 --- a/crates/hipfire-arch-qwen2/src/qwen2.rs +++ b/crates/hipfire-arch-qwen2/src/qwen2.rs @@ -276,26 +276,29 @@ impl Qwen2Weights { /// Release every GPU buffer back to the pool. Consumes self. /// Mirrors `LlamaWeights::free_gpu` and `Qwen35Weights::free_gpu` - /// — the daemon calls this on unload to actually return VRAM. + /// — the daemon calls this on unload to actually return VRAM. Linear + /// weights go through `WeightTensor::free_all` so any PARO / AWQ sidecar + /// (none are allocated for Qwen2 today, but the forward already reads + /// `awq_scale`) is released with the buffer instead of leaking per reload. pub fn free_gpu(self, gpu: &mut Gpu) { let _ = gpu.free_tensor(self.token_embd); let _ = gpu.free_tensor(self.output_norm); if !self.tied_lm_head { - let _ = gpu.free_tensor(self.output.buf); + self.output.free_all(gpu); } for l in self.layers { let _ = gpu.free_tensor(l.attn_norm); - let _ = gpu.free_tensor(l.wq.buf); + l.wq.free_all(gpu); let _ = gpu.free_tensor(l.wq_bias); - let _ = gpu.free_tensor(l.wk.buf); + l.wk.free_all(gpu); let _ = gpu.free_tensor(l.wk_bias); - let _ = gpu.free_tensor(l.wv.buf); + l.wv.free_all(gpu); let _ = gpu.free_tensor(l.wv_bias); - let _ = gpu.free_tensor(l.wo.buf); + l.wo.free_all(gpu); let _ = gpu.free_tensor(l.ffn_norm); - let _ = gpu.free_tensor(l.w_gate.buf); - let _ = gpu.free_tensor(l.w_up.buf); - let _ = gpu.free_tensor(l.w_down.buf); + l.w_gate.free_all(gpu); + l.w_up.free_all(gpu); + l.w_down.free_all(gpu); } } } diff --git a/crates/hipfire-arch-qwen35/map.md b/crates/hipfire-arch-qwen35/map.md index 94a8f3efd..ea268919d 100644 --- a/crates/hipfire-arch-qwen35/map.md +++ b/crates/hipfire-arch-qwen35/map.md @@ -42,7 +42,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | [`src/qwen35/ep_batch.rs`](src/qwen35/ep_batch.rs) | 4,800 | 20 | 7 | | [`src/qwen35/forward.rs`](src/qwen35/forward.rs) | 6,255 | 31 | 12 | | [`src/qwen35/load.rs`](src/qwen35/load.rs) | 4,906 | 10 | 0 | -| [`src/qwen35/prefill.rs`](src/qwen35/prefill.rs) | 9,312 | 11 | 48 | +| [`src/qwen35/prefill.rs`](src/qwen35/prefill.rs) | 9,381 | 11 | 49 | | [`src/qwen35/weights.rs`](src/qwen35/weights.rs) | 1,971 | 43 | 10 | | [`src/qwen35.rs`](src/qwen35.rs) | 63 | 7 | 0 | | [`src/scheduler.rs`](src/scheduler.rs) | 142 | 3 | 4 | @@ -97,6 +97,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 29 modules · 57,281 lines · 436 public items · 189 tests · 4 examples +- 29 modules · 57,350 lines · 436 public items · 190 tests · 4 examples diff --git a/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs b/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs index 418e23b44..9c65fb3b1 100644 --- a/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs +++ b/crates/hipfire-arch-qwen35/src/qwen35/prefill.rs @@ -1543,17 +1543,16 @@ pub(crate) fn is_batchable_la(dt: DType, arch: &str) -> bool { // (gfx1100/1101/1102/1150/1151 + gfx1200/1201) but gate the gfx11 half // behind HIPFIRE_MQV2_GFX11_WMMA != "0" — setting // HIPFIRE_MQV2_GFX11_WMMA=0 restores the per-token fallback ONLY on - // gfx11, leaving gfx12 untouched. Lockstep with the HasWmma predicate - // on GemmMq*G256V2* keys and with gemm_mq*g256v2's has_wmma() guard. + // gfx11, leaving gfx12 untouched. Delegates to the shared + // `hipfire_runtime::llama::mqv2_wmma_batchable` rule (shared home for + // the dtype/arch/kill-switch set). NOTE: `llama::is_batchable_la` does + // NOT delegate to it — the llama chunk path has no V2 arms, so llama + // refuses V2 everywhere; only this qwen35 caller admits V2. Lockstep with + // the HasWmma predicate on GemmMq*G256V2* keys and with + // gemm_mq*g256v2's has_wmma() guard. // MQ4CG256 (qt45) remains gfx12-only until its gfx11 sibling lands. - let mqv2_with_wmma = matches!( + let mqv2_with_wmma = llama::mqv2_wmma_batchable( dt, - DType::MQ4G256V2 - | DType::MQ6G256V2 - | DType::MQ5G256V2 - | DType::MQ3G256V2 - | DType::MQ2G256V2 - ) && mqv2_gfx11_wmma_enabled_from_env( hipfire_config::developer_var("HIPFIRE_MQV2_GFX11_WMMA") .ok() .as_deref(), @@ -1582,24 +1581,6 @@ pub(crate) fn is_batchable_la(dt: DType, arch: &str) -> bool { || bf16_with_gfx942 } -/// Helper for MQ2/3/4/5/6G256V2 (qt44,47-50) batched prefill admit: gfx12 always, gfx11 -/// gated by HIPFIRE_MQV2_GFX11_WMMA != "0". Public for testability, mirrors -/// `mq6_batched_admit_enabled_from_env` / `q8_prefill_wmma_enabled_from_env`. -/// `value` is the raw env var (None = unset → default ON); only Some("0") -/// disables the gfx11 path. Gfx12 is unaffected by the env var. -pub(crate) fn mqv2_gfx11_wmma_enabled_from_env(value: Option<&str>, arch: &str) -> bool { - let gfx11_enabled = value != Some("0"); - if matches!(arch, "gfx1200" | "gfx1201") { - true - } else if matches!( - arch, - "gfx1100" | "gfx1101" | "gfx1102" | "gfx1150" | "gfx1151" - ) { - gfx11_enabled - } else { - false - } -} /// Single source of truth for per-layer batchability and checked geometry. /// Called by `validate_ep_batch_compatibility`, `prefill_batch_pbs_eligible`, /// `fa_batched_ok` guard, and later EP state preflight. Validates every @@ -8152,44 +8133,44 @@ mod tests { #[test] fn qwen35_is_batchable_la_mq4_v2_env_escape() { // HIPFIRE_MQV2_GFX11_WMMA=0 restores fallback ONLY on gfx11; gfx12 - // remains admitted. Use the helper directly to avoid global env + // remains admitted. Use the shared helper directly to avoid global env // mutation flakiness in parallel tests — is_batchable_la delegates - // to this helper verbatim. + // to `llama::mqv2_wmma_batchable`, which calls this helper verbatim. for arch in ["gfx1100", "gfx1101", "gfx1102", "gfx1150", "gfx1151"] { assert!( - !mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch), + !llama::mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch), "env=0 should disable {arch}" ); assert!( - mqv2_gfx11_wmma_enabled_from_env(None, arch), + llama::mqv2_gfx11_wmma_enabled_from_env(None, arch), "unset should enable {arch}" ); assert!( - mqv2_gfx11_wmma_enabled_from_env(Some("1"), arch), + llama::mqv2_gfx11_wmma_enabled_from_env(Some("1"), arch), "env=1 should enable {arch}" ); } for arch in ["gfx1200", "gfx1201"] { assert!( - mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch), + llama::mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch), "gfx12 unaffected by env=0 on {arch}" ); assert!( - mqv2_gfx11_wmma_enabled_from_env(None, arch), + llama::mqv2_gfx11_wmma_enabled_from_env(None, arch), "gfx12 enabled without env on {arch}" ); } for arch in ["gfx1010", "gfx942", "gfx1030", "gfx1103", "gfx1152"] { assert!( - !mqv2_gfx11_wmma_enabled_from_env(None, arch), + !llama::mqv2_gfx11_wmma_enabled_from_env(None, arch), "non-WMMA {arch} must never admit" ); assert!( - !mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch), + !llama::mqv2_gfx11_wmma_enabled_from_env(Some("0"), arch), "non-WMMA {arch} with env=0" ); assert!( - !mqv2_gfx11_wmma_enabled_from_env(Some("1"), arch), + !llama::mqv2_gfx11_wmma_enabled_from_env(Some("1"), arch), "non-WMMA {arch} with env=1" ); } @@ -8241,6 +8222,94 @@ mod tests { assert_eq!(rdna_compute::MQ4V2_GROUP_BYTES, 136); } + #[test] + fn mqv2_admit_llama_qwen35_lockstep() { + // True contract (PR #690 hw-gate regression): llama and qwen35 agree + // on every NON-V2 dtype, but for the V2 family they deliberately + // diverge — qwen35's `forward_prefill_chunk` has V2 dispatch arms + // (206 hits) so it admits V2 via the shared + // `llama::mqv2_wmma_batchable` rule, while llama's chunk path has no + // V2 arms (`qkv_is_mq`/`wo_is_mq`/`ffn_is_mq`/`w_down_is_mq` list + // only V1 dtypes) so `llama::is_batchable_la` refuses V2 everywhere + // and stays on per-token decode. Admitting V2 to the llama path + // would skip the FWHT rotate and run V1 `hfq4g256` launchers on V2 + // blobs — silently incoherent prefill. + // Non-V2 agreement across the 5-arch sample. + let non_v2 = [ + DType::MQ4G256, + DType::HFQ4G256, + DType::MQ6G256, + DType::MQ3G256, + DType::MFP4G32, + DType::Q8_0, + ]; + for dt in non_v2 { + for arch in ["gfx1100", "gfx1151", "gfx1201", "gfx1030", "gfx1010"] { + assert_eq!( + llama::is_batchable_la(dt, arch), + is_batchable_la(dt, arch), + "lockstep drift for {dt:?} on {arch}" + ); + } + } + // V2 divergence: qwen35 admits on gfx11/gfx12 (kill-switch at its + // default ON here — both gates read `HIPFIRE_MQV2_GFX11_WMMA` + // identically, so with the var unset gfx11 admits), refuses + // pre-WMMA; llama refuses on all 5 arches. + let v2 = [ + DType::MQ4G256V2, + DType::MQ6G256V2, + DType::MQ5G256V2, + DType::MQ3G256V2, + DType::MQ2G256V2, + DType::MQ4CG256, + ]; + for dt in v2 { + for arch in ["gfx1100", "gfx1151"] { + // MQ4CG256 is gfx12-only by intent in BOTH callers. + if dt == DType::MQ4CG256 { + assert!( + !is_batchable_la(dt, arch), + "qwen35 must refuse {dt:?} on {arch}" + ); + } else { + assert!( + is_batchable_la(dt, arch), + "qwen35 should admit {dt:?} on {arch}" + ); + } + assert!( + !llama::is_batchable_la(dt, arch), + "llama must refuse {dt:?} on {arch}" + ); + } + assert!( + is_batchable_la(dt, "gfx1201"), + "qwen35 should admit {dt:?} on gfx1201" + ); + assert!( + !llama::is_batchable_la(dt, "gfx1201"), + "llama must refuse {dt:?} on gfx1201" + ); + for arch in ["gfx1030", "gfx1010"] { + assert!( + !is_batchable_la(dt, arch), + "qwen35 must refuse {dt:?} on {arch}" + ); + assert!( + !llama::is_batchable_la(dt, arch), + "llama must refuse {dt:?} on {arch}" + ); + } + } + // Absolute pins so the test also fails if the shared rule itself + // regresses, not just on caller drift. + assert!(is_batchable_la(DType::MQ4G256V2, "gfx1201")); + assert!(!is_batchable_la(DType::MQ4G256V2, "gfx1030")); + assert!(!is_batchable_la(DType::MQ4CG256, "gfx1100")); + assert!(!llama::is_batchable_la(DType::MQ4G256V2, "gfx1201")); + } + #[test] fn qwen35_v2_dense_keys_are_exact_no_hfq4_default() { // Contract: every admitted V2 dtype maps 1:1 to its exact V2 kernel diff --git a/crates/hipfire-cli/Cargo.toml b/crates/hipfire-cli/Cargo.toml index bbe77622f..b82cf94d3 100644 --- a/crates/hipfire-cli/Cargo.toml +++ b/crates/hipfire-cli/Cargo.toml @@ -19,6 +19,7 @@ http-body-util = "0.1" hyper = { version = "1", features = ["http1", "server"] } hyper-util = { version = "0.1", features = ["tokio"] } libc = "0.2" +md5 = "0.8" hipfire-config = { path = "../hipfire-config" } hipfire-client = { path = "../hipfire-client" } hipfire-registry = { path = "../hipfire-registry" } diff --git a/crates/hipfire-cli/map.md b/crates/hipfire-cli/map.md index d87967dec..89475e63a 100644 --- a/crates/hipfire-cli/map.md +++ b/crates/hipfire-cli/map.md @@ -23,11 +23,11 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| | [`src/bench_concurrency.rs`](src/bench_concurrency.rs) | 720 | 21 | 9 | -| [`src/main.rs`](src/main.rs) | 9,692 | 0 | 65 | +| [`src/main.rs`](src/main.rs) | 10,762 | 0 | 87 | | [`src/serve/complete.rs`](src/serve/complete.rs) | 6,754 | 0 | 89 | | [`src/serve/http.rs`](src/serve/http.rs) | 1,089 | 0 | 6 | | [`src/serve/metrics.rs`](src/serve/metrics.rs) | 328 | 0 | 5 | -| [`src/serve/mod.rs`](src/serve/mod.rs) | 2,116 | 2 | 16 | +| [`src/serve/mod.rs`](src/serve/mod.rs) | 2,127 | 2 | 16 | | [`src/setup.rs`](src/setup.rs) | 1,529 | 0 | 16 | ### Public API surface @@ -43,7 +43,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Dependencies (from `Cargo.toml`) - path: `hipfire-client`, `hipfire-config`, `hipfire-registry`, `hipfire-runtime`, `saddle-core` -- external: `anyhow`, `bytes`, `clap`, `ctrlc`, `http-body-util`, `hyper`, `hyper-util`, `libc`, `serde`, `serde_json`, `sha2`, `tokio`, `tokio-util`, `ureq` +- external: `anyhow`, `bytes`, `clap`, `ctrlc`, `http-body-util`, `hyper`, `hyper-util`, `libc`, `md5`, `serde`, `serde_json`, `sha2`, `tokio`, `tokio-util`, `ureq` - dev: — - build: — @@ -53,6 +53,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 7 modules · 22,228 lines · 23 public items · 206 tests · 0 examples +- 7 modules · 23,309 lines · 23 public items · 228 tests · 0 examples diff --git a/crates/hipfire-cli/src/main.rs b/crates/hipfire-cli/src/main.rs index 86764ac22..40b00f8ad 100644 --- a/crates/hipfire-cli/src/main.rs +++ b/crates/hipfire-cli/src/main.rs @@ -489,6 +489,10 @@ pub(crate) struct BenchArgs { /// Prompt words for the standard benchmark. #[arg(num_args = 0..)] prompt: Vec, + /// Read the standard-benchmark prompt verbatim from a file (raw bytes, + /// no trimming). Mutually exclusive with positional PROMPT words. + #[arg(long, conflicts_with = "prompt")] + prompt_file: Option, } #[derive(Args, Debug)] @@ -1618,6 +1622,7 @@ pub(crate) fn pull_command(paths: &Paths, args: PullArgs) -> Result<()> { ("TriAttention", entry.triattn.as_ref()), ("MTP", entry.mtp.as_ref()), ("DSpark", entry.dspark.as_ref()), + ("DFlash", entry.dflash.as_ref()), ] { let Some(sidecar) = sidecar else { continue; @@ -1770,14 +1775,28 @@ fn report_progress(downloaded: u64, total: Option, elapsed: Duration) { fn rm_command(paths: &Paths, args: RmArgs) -> Result<()> { let loaded = load_registry(&paths.registry); - let resolved = loaded.registry.model(&args.model); - let path = find_model_path(paths, &loaded.registry, &args.model) + rm_with_registry(paths, &loaded.registry, args) +} + +/// Remove one model and its sidecars. A declared DFlash draft sidecar is +/// shared: several registry entries can name the same `dflash.file` (e.g. +/// `qwen3.8:27b`, `qwen3.8:27b-mq4-pro`, and `qwen3.8:27b-mq4-xt` all declare +/// `qwen38-27b-dflash-mq4.hfq`). Deleting it while a sibling declarer is +/// still on disk leaves those siblings running AR under `dflash_mode=auto` +/// or refusing to load under `on`, so the sidecar is kept — with one stderr +/// line — whenever any OTHER entry declaring the same file still has its own +/// target file present in the models dir. +fn rm_with_registry(paths: &Paths, registry: &RegistryV1, args: RmArgs) -> Result<()> { + let resolved = registry_entry_for_path(paths, registry, &args.model); + let path = find_model_path(paths, registry, &args.model) .unwrap_or_else(|| paths.models.join(&args.model)); if !path.is_file() { bail!("model not found: {}", path.display()); } let mut targets = BTreeSet::from([path.clone()]); - if let Some((_, entry)) = resolved { + // A shared DFlash sidecar that must survive this removal: (file, keepers). + let mut kept_sidecar: Option<(String, String)> = None; + if let Some((tag, entry)) = resolved { targets.extend( [&entry.triattn, &entry.mtp, &entry.dspark] .into_iter() @@ -1785,6 +1804,31 @@ fn rm_command(paths: &Paths, args: RmArgs) -> Result<()> { .map(|sidecar| paths.models.join(&sidecar.file)) .filter(|path| path.is_file()), ); + if let Some(sidecar) = entry.dflash.as_ref() { + let sidecar_path = paths.models.join(&sidecar.file); + if sidecar_path.is_file() { + // `models` is a BTreeMap, so keepers list in sorted tag order. + let keepers: Vec<&str> = registry + .models + .iter() + .filter(|(other_tag, other)| { + other_tag.as_str() != tag + && other.file != entry.file + && other + .dflash + .as_ref() + .is_some_and(|other_sidecar| other_sidecar.file == sidecar.file) + && paths.models.join(&other.file).is_file() + }) + .map(|(other_tag, _)| other_tag.as_str()) + .collect(); + if keepers.is_empty() { + targets.insert(sidecar_path); + } else { + kept_sidecar = Some((sidecar.file.clone(), keepers.join(", "))); + } + } + } } if let (Some(parent), Some(file)) = ( path.parent(), @@ -1828,14 +1872,16 @@ fn rm_command(paths: &Paths, args: RmArgs) -> Result<()> { .with_context(|| format!("failed to remove {}", target.display()))?; println!("removed {}", target.display()); } + if let Some((file, keepers)) = kept_sidecar { + eprintln!("keeping DFlash sidecar {file}: still declared by {keepers}"); + } Ok(()) } fn run_command(paths: &Paths, args: RunArgs) -> Result<()> { let loaded_registry = load_registry(&paths.registry); let registry = &loaded_registry.registry; - let (canonical, entry) = registry - .model(&args.model) + let (canonical, entry) = registry_entry_for_path(paths, registry, &args.model) .map(|(tag, entry)| (Some(tag.to_owned()), Some(entry))) .unwrap_or((None, None)); let mut model_path = find_model_path(paths, registry, &args.model); @@ -1944,10 +1990,13 @@ fn run_command(paths: &Paths, args: RunArgs) -> Result<()> { let mut params = load_params( &resolved, entry, + &paths.models, &model_path, max_tokens, args.kv_mode.as_deref(), args.kv_backend.as_deref(), + canonical.as_deref(), + args.model_draft.is_some(), )?; let selector = args .speculation @@ -1963,6 +2012,16 @@ fn run_command(paths: &Paths, args: RunArgs) -> Result<()> { apply_speculation_selector(&mut params, "dflash")?; } } + // Registry sidecar for a final auto/on selector the config-time + // load_params could not see (config-off + `run --spec dflash`); a + // no-op when load_params already resolved or an explicit draft won. + resolve_dflash_sidecar( + &mut params, + entry, + &paths.models, + &model_path, + canonical.as_deref(), + )?; if let Some(window) = args.draft_max { if !(1..=32).contains(&window) { bail!("--draft-max must be between 1 and 32"); @@ -2395,6 +2454,36 @@ fn scan_local_models(local: &[PathBuf], search: &str, mode: MatchMode) -> Vec( + paths: &Paths, + registry: &'registry RegistryV1, + input: &str, +) -> Option<(&'registry str, &'registry ModelEntry)> { + let candidate = Path::new(input); + if input.contains('/') || input.contains('\\') || candidate.is_file() { + let canonical_input = fs::canonicalize(candidate).ok()?; + return registry.models.iter().find_map(|(tag, entry)| { + let canonical_installed = fs::canonicalize(paths.models.join(&entry.file)).ok()?; + (canonical_installed == canonical_input).then(|| (tag.as_str(), entry)) + }); + } + registry.model(input) +} + pub(crate) fn find_model_path( paths: &Paths, registry: &RegistryV1, @@ -2488,10 +2577,13 @@ pub(crate) fn find_model_path( pub(crate) fn load_params( resolved: &hipfire_config::ResolvedConfig, entry: Option<&ModelEntry>, + models_dir: &Path, model_path: &Path, max_tokens: u64, kv_override: Option<&str>, kv_backend_override: Option<&str>, + tag: Option<&str>, + explicit_draft: bool, ) -> Result { let configured_max_seq = config_u64(resolved, "memory.max_seq")?; let max_seq = configured_max_seq.max(max_tokens.saturating_add(1024)); @@ -2573,9 +2665,78 @@ pub(crate) fn load_params( let selector = config_string(resolved, "speculation.mode")?; apply_speculation_selector(&mut params, &selector)?; project_dflash_draft(&mut params, developer_dflash_draft(resolved)); + if !explicit_draft { + // A CLI `--model-draft` (projected by the caller after this returns) + // always wins, so skip sidecar resolution — and its `on` fail-closed + // bail — when one was given. + resolve_dflash_sidecar(&mut params, entry, models_dir, model_path, tag)?; + } Ok(params) } +/// Resolve a registry-declared DFlash sidecar into `params["draft"]`. +/// +/// Call only once the final `dflash_mode` is known. When the mode is `auto` +/// or `on`, no explicit draft is set (`params["draft"]`, e.g. from +/// `developer.dflash_draft`), and `entry.dflash` names a pulled file, wire +/// it: `on` without the file fails closed, `auto` logs one line and runs AR. +/// With no entry at all the artifact is not registry-managed (e.g. a path +/// that merely shares a basename with an entry file): `auto` runs AR as a +/// bare artifact, but `on` fails closed instead of silently running AR. +/// The sidecar is looked up in `models_dir` first — `find_model_path` +/// canonicalizes, so a symlinked target's parent is wherever the artifact +/// really lives, not the models directory the draft was pulled into — then +/// next to the target. An explicit draft always wins; a final `off` never +/// carries a draft (`project_dflash_draft` strips it) and returns early. +fn resolve_dflash_sidecar( + params: &mut serde_json::Value, + entry: Option<&ModelEntry>, + models_dir: &Path, + model_path: &Path, + tag: Option<&str>, +) -> Result<()> { + if !matches!(params["dflash_mode"].as_str(), Some("auto" | "on")) { + return Ok(()); + } + if params + .get("draft") + .and_then(serde_json::Value::as_str) + .is_some_and(|draft| !draft.is_empty()) + { + return Ok(()); + } + let Some(sidecar) = entry.and_then(|entry| entry.dflash.as_ref()) else { + if params["dflash_mode"].as_str() == Some("on") && model_path.is_file() { + bail!( + "DFlash draft required (dflash_mode=on) but {} is not a registry-managed artifact; pass developer.dflash_draft or use the registry tag", + model_path.display() + ); + } + return Ok(()); + }; + let beside_target = model_path.parent().unwrap_or_else(|| Path::new(".")); + let candidate = [models_dir, beside_target] + .into_iter() + .map(|dir| dir.join(&sidecar.file)) + .find(|candidate| candidate.is_file()); + if let Some(candidate) = candidate { + params["draft"] = serde_json::json!(candidate.display().to_string()); + return Ok(()); + } + let tag = tag.unwrap_or(""); + if params["dflash_mode"].as_str() == Some("on") { + bail!( + "DFlash draft {} is not pulled; run `hipfire pull {tag}` or set developer.dflash_draft", + sidecar.file + ); + } + eprintln!( + "[hipfire] DFlash draft {} not pulled; running AR — `hipfire pull {tag}`", + sidecar.file + ); + Ok(()) +} + /// Project snapshotted `developer.dflash_draft` after the effective speculation selector. /// /// Call only once final `dflash_mode` is known. Config-off must not carry a draft; @@ -3615,6 +3776,69 @@ fn sample_stats(values: &[f64]) -> Option { }) } +/// Default standard-bench prompt. Historical numbers depend on its exact +/// bytes; do not change it. +const BENCH_DEFAULT_PROMPT: &str = "Explain the theory of general relativity in simple terms."; + +/// Below this tokenized prompt length the measured `prefill_tok_s` is launch +/// overhead, not prefill throughput (363 tok/s at ~24 tokens vs 886 at 4.4k +/// on a 7900 XTX with the same binary). +const BENCH_PREFILL_EVIDENCE_TOKENS: u64 = 256; + +/// Resolve the standard-benchmark prompt: `--prompt-file` reads the file +/// verbatim (raw bytes, no trimming — one newline can move τ by 17%), else +/// the positional words joined with spaces, else the historical default. +/// The clap `conflicts_with` on `--prompt-file` covers CLI parsing; the +/// explicit check here covers programmatically built args. +fn resolve_bench_prompt(args: &BenchArgs) -> Result { + if let Some(path) = args.prompt_file.as_deref() { + if !args.prompt.is_empty() { + bail!("--prompt-file cannot be combined with a positional prompt"); + } + let bytes = fs::read(path) + .with_context(|| format!("failed to read --prompt-file {}", path.display()))?; + return String::from_utf8(bytes) + .with_context(|| format!("--prompt-file {} is not valid UTF-8", path.display())); + } + Ok(if args.prompt.is_empty() { + BENCH_DEFAULT_PROMPT.to_owned() + } else { + args.prompt.join(" ") + }) +} + +/// Hex md5 of the exact prompt bytes sent, so two bench numbers are only +/// compared when their prompts are byte-identical. +fn bench_prompt_md5(prompt: &str) -> String { + format!("{:x}", md5::compute(prompt.as_bytes())) +} + +/// Short-prompt caveat: below 256 prompt tokens `prefill_tok_s` measures +/// launch overhead, not prefill throughput. +fn bench_prompt_warning(prompt_tokens: u64) -> Option { + (prompt_tokens < BENCH_PREFILL_EVIDENCE_TOKENS).then(|| { + format!( + "prompt is {prompt_tokens} tokens; prefill_tok_s at this length measures launch overhead, not prefill throughput — use --prompt-file with ≥256 tokens for a prefill number" + ) + }) +} + +/// Prompt length as the daemon reports it on the `done` event: +/// `prefill_tokens` (rows actually prefilled) plus `cached_tokens` (prefix +/// served from the prompt cache). `None` when the event carries neither. +fn bench_prompt_tokens_from_done(done: &serde_json::Value) -> Option { + let prefill = done + .get("prefill_tokens") + .and_then(serde_json::Value::as_u64); + let cached = done + .get("cached_tokens") + .and_then(serde_json::Value::as_u64); + match (prefill, cached) { + (None, None) => None, + (p, c) => Some(p.unwrap_or(0) + c.unwrap_or(0)), + } +} + fn bench_command(paths: &Paths, args: BenchArgs) -> Result<()> { if args.runs == 0 { bail!("--runs must be positive"); @@ -3665,11 +3889,9 @@ fn bench_command(paths: &Paths, args: BenchArgs) -> Result<()> { return bench_experimental(paths, &args); } let (mut engine, loaded, pre_diag, post_diag) = open_bench_engine(paths, &args, None)?; - let prompt = if args.prompt.is_empty() { - "Explain the theory of general relativity in simple terms.".to_owned() - } else { - args.prompt.join(" ") - }; + let prompt = resolve_bench_prompt(&args)?; + let prompt_md5 = bench_prompt_md5(&prompt); + let prompt_chars = prompt.chars().count() as u64; eprintln!("hipfire bench"); eprintln!(" model: {}", args.model); eprintln!( @@ -3688,6 +3910,8 @@ fn bench_command(paths: &Paths, args: BenchArgs) -> Result<()> { ); eprintln!(" runs: {}", args.runs); eprintln!(" max_tokens: {}", args.max_tokens); + eprintln!(" prompt_md5: {prompt_md5}"); + eprintln!(" prompt_chars: {prompt_chars}"); if args.matrix || args.redline { bench_matrix(&mut engine, &args, &loaded, &post_diag) } else { @@ -3700,6 +3924,7 @@ fn bench_command(paths: &Paths, args: BenchArgs) -> Result<()> { let mut prefill = Vec::new(); let mut wall = Vec::new(); let mut ttft = Vec::new(); + let mut prompt_tokens: Option = None; for _ in 0..args.runs { let done = bench_generate_with_reasoning( &mut engine, @@ -3707,6 +3932,14 @@ fn bench_command(paths: &Paths, args: BenchArgs) -> Result<()> { args.max_tokens as u64, args.reasoning_on, )?; + // Every run uses the same prompt, so the daemon's tokenized + // prompt length is run-invariant; keep the first report. The + // done event reports the prompt as `prefill_tokens` (rows the + // engine actually prefilled) plus `cached_tokens` (prefix served + // from the prompt cache); the prompt is their sum. + if prompt_tokens.is_none() { + prompt_tokens = bench_prompt_tokens_from_done(&done); + } if let Some(value) = done.get("decode_tok_s").and_then(serde_json::Value::as_f64) { decode.push(value); } @@ -3726,6 +3959,19 @@ fn bench_command(paths: &Paths, args: BenchArgs) -> Result<()> { std::io::stderr().flush()?; } eprintln!(); + // Below 256 prompt tokens prefill_tok_s is launch overhead, not + // prefill throughput — say so in the report instead of leaving a + // bare number that invites a wrong comparison. + let warnings: Vec = prompt_tokens + .and_then(bench_prompt_warning) + .into_iter() + .collect(); + if let Some(tokens) = prompt_tokens { + eprintln!(" prompt_tokens: {tokens}"); + } + for warning in &warnings { + eprintln!(" warning: {warning}"); + } let report = serde_json::json!({ "protocol": "native-generate-v1", "model": args.model, @@ -3735,6 +3981,10 @@ fn bench_command(paths: &Paths, args: BenchArgs) -> Result<()> { "max_tokens": args.max_tokens, "runs": args.runs, "batch": 1, + "prompt_tokens": prompt_tokens, + "prompt_md5": prompt_md5, + "prompt_chars": prompt_chars, + "warnings": warnings, "decode_tok_s": sample_stats(&decode), "prefill_tok_s": sample_stats(&prefill), "wall_tok_s": sample_stats(&wall), @@ -3940,8 +4190,7 @@ fn open_bench_engine( serde_json::Value, )> { let registry = load_registry(&paths.registry).registry; - let (tag, entry) = registry - .model(&args.model) + let (tag, entry) = registry_entry_for_path(paths, ®istry, &args.model) .map(|(tag, entry)| (Some(tag.to_owned()), Some(entry.clone()))) .unwrap_or((None, None)); let mut path = find_model_path(paths, ®istry, &args.model); @@ -3991,14 +4240,26 @@ fn open_bench_engine( let mut params = load_params( &resolved, entry.as_ref(), + &paths.models, &path, max_tokens, args.kv_mode.as_deref(), args.kv_backend.as_deref(), + tag.as_deref(), + false, )?; if let Some(selector) = args.speculation.as_deref() { apply_speculation_selector(&mut params, selector)?; } + // Registry sidecar for a final auto/on selector the config-time + // load_params could not see (config-off + `bench --spec dflash`). + resolve_dflash_sidecar( + &mut params, + entry.as_ref(), + &paths.models, + &path, + tag.as_deref(), + )?; if args.matrix || args.redline { let requested = longest_prefill.max(longest_decode).saturating_add(32); let configured = params["max_seq"].as_u64().unwrap_or(0); @@ -4225,11 +4486,7 @@ fn bench_experimental(paths: &Paths, args: &BenchArgs) -> Result<()> { bail!("--exp requires RDNA2 (gfx1030/gfx1031), detected {arch}"); } let _ = bench_generate(&mut engine, "Hello", 16)?; - let prompt = if args.prompt.is_empty() { - "Explain the theory of general relativity in simple terms.".to_owned() - } else { - args.prompt.join(" ") - }; + let prompt = resolve_bench_prompt(args)?; let mut samples = Vec::new(); for _ in 0..args.runs { let done = bench_generate(&mut engine, &prompt, 128)?; @@ -4285,6 +4542,7 @@ fn profile_command(paths: &Paths, args: ProfileArgs) -> Result<()> { backend: "both".to_owned(), workload: "both".to_owned(), prompt: Vec::new(), + prompt_file: None, }; let (mut engine, _, _, _) = open_bench_engine(paths, &bench, None)?; let _ = bench_generate(&mut engine, "Hello", 1)?; @@ -6254,7 +6512,18 @@ mod tests { fs::write(&sidecar_path, b"sidecar").unwrap(); let defaults = resolve(Vec::::new()).unwrap(); - let params = load_params(&defaults, Some(entry), &model_path, 64, None, None).unwrap(); + let params = load_params( + &defaults, + Some(entry), + &model_path.parent().unwrap(), + &model_path, + 64, + None, + None, + None, + false, + ) + .unwrap(); assert_eq!(params["cask"], false); assert_eq!(params["cask_handoff_tokens"], 0); assert_eq!(params["cask_sidecar"], ""); @@ -6269,7 +6538,18 @@ mod tests { layer: explicit, }]) .unwrap(); - let params = load_params(&enabled, Some(entry), &model_path, 64, None, None).unwrap(); + let params = load_params( + &enabled, + Some(entry), + &model_path.parent().unwrap(), + &model_path, + 64, + None, + None, + None, + false, + ) + .unwrap(); assert_eq!(params["cask"], false); assert_eq!(params["cask_sidecar"], sidecar_path.display().to_string()); assert_eq!(params["prefill_compression"], "off"); @@ -6280,8 +6560,18 @@ mod tests { pub(crate) fn load_params_forwards_explicit_vmm_backend() { let defaults = resolve(Vec::::new()).unwrap(); let model_path = PathBuf::from("/tmp/test-model.mq4"); - let params = - load_params(&defaults, None, &model_path, 64, Some("q8"), Some("vmm")).unwrap(); + let params = load_params( + &defaults, + None, + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + Some("vmm"), + None, + false, + ) + .unwrap(); assert_eq!(params["kv_backend"], "vmm"); } @@ -6289,7 +6579,18 @@ mod tests { pub(crate) fn load_params_defaults_to_schema_contiguous_backend() { let defaults = resolve(Vec::::new()).unwrap(); let model_path = PathBuf::from("/tmp/test-model.mq4"); - let params = load_params(&defaults, None, &model_path, 64, Some("q8"), None).unwrap(); + let params = load_params( + &defaults, + None, + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + None, + false, + ) + .unwrap(); assert_eq!(params["kv_backend"], "contiguous"); assert_eq!(params["max_seq"], 32768); } @@ -6587,16 +6888,29 @@ mod tests { let params = load_params( &resolved, Some(entry), + &model_path.parent().unwrap(), &model_path, 64, Some("q8"), Some("contiguous"), + None, + false, ) .unwrap(); assert_eq!(params["kv_backend"], "contiguous"); // Without explicit override, load_params uses the resolved vmm. - let params2 = - load_params(&resolved, Some(entry), &model_path, 64, Some("q8"), None).unwrap(); + let params2 = load_params( + &resolved, + Some(entry), + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + None, + false, + ) + .unwrap(); assert_eq!(params2["kv_backend"], "vmm"); assert_eq!(params2["max_seq"], 262144); @@ -6695,7 +7009,18 @@ mod tests { pub(crate) fn load_params_only_forwards_explicit_deepseek4_expert_fanout() { let model_path = PathBuf::from("/tmp/test-model.mq2r"); let defaults = resolve(Vec::::new()).unwrap(); - let params = load_params(&defaults, None, &model_path, 64, Some("q8"), None).unwrap(); + let params = load_params( + &defaults, + None, + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + None, + false, + ) + .unwrap(); assert_eq!(params["deepseek4_compute_placement"], "single"); assert!(params.get("deepseek4_experts_per_token").is_none()); @@ -6710,7 +7035,18 @@ mod tests { layer: explicit, }]) .unwrap(); - let params = load_params(&resolved, None, &model_path, 64, Some("q8"), None).unwrap(); + let params = load_params( + &resolved, + None, + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + None, + false, + ) + .unwrap(); assert_eq!(params["deepseek4_experts_per_token"], 4); } @@ -6731,10 +7067,13 @@ mod tests { let params = load_params( &resolved, None, + Path::new("/tmp/test-model.mq2r").parent().unwrap(), Path::new("/tmp/test-model.mq2r"), 64, Some("q8"), None, + None, + false, ) .unwrap(); assert_eq!(params["deepseek4_compute_placement"], raw); @@ -6756,10 +7095,615 @@ mod tests { .unwrap(); let model_path = PathBuf::from("/tmp/test-model.mq4"); - let params = load_params(&resolved, None, &model_path, 64, Some("q8"), None).unwrap(); + let params = load_params( + &resolved, + None, + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + None, + false, + ) + .unwrap(); assert_eq!(params["draft"], draft); } + fn dflash_sidecar_entry(draft_file: &str) -> ModelEntry { + ModelEntry { + repo: "hipfire-models/qwen3.5-9b".into(), + file: "qwen3.5-9b.mq4".into(), + size_gb: 5.31, + min_vram_gb: 6.8, + desc: "test target".into(), + dflash: Some(hipfire_registry::Sidecar { + file: draft_file.into(), + sha256: None, + size_bytes: None, + }), + ..Default::default() + } + } + + fn resolved_with_dflash_mode( + mode: &str, + draft: Option<&str>, + ) -> hipfire_config::ResolvedConfig { + let mut explicit = ConfigLayer::default(); + explicit.set_cli("speculation.dflash", mode).unwrap(); + if let Some(draft) = draft { + explicit.set_cli("developer.dflash_draft", draft).unwrap(); + } + resolve([NamedLayer { + source: ConfigSource::OneShot { + argument: format!("speculation.dflash={mode}"), + }, + layer: explicit, + }]) + .unwrap() + } + + #[test] + pub(crate) fn load_params_resolves_registry_dflash_sidecar_when_present() { + // (b) auto + pulled draft file → params["draft"] points at it. + let paths = test_paths("dflash-sidecar-present"); + fs::create_dir_all(&paths.models).unwrap(); + let model_path = paths.models.join("qwen3.5-9b.mq4"); + fs::write(&model_path, b"model").unwrap(); + let draft_path = paths.models.join("qwen35-9b-dflash-mq4.hfq"); + fs::write(&draft_path, b"draft").unwrap(); + let entry = dflash_sidecar_entry("qwen35-9b-dflash-mq4.hfq"); + let resolved = resolved_with_dflash_mode("auto", None); + let params = load_params( + &resolved, + Some(&entry), + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + Some("qwen3.5:9b"), + false, + ) + .unwrap(); + assert_eq!(params["dflash_mode"], "auto"); + assert_eq!(params["draft"], draft_path.display().to_string()); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[cfg(unix)] + #[test] + pub(crate) fn load_params_finds_sidecar_in_models_dir_for_symlinked_target() { + // find_model_path canonicalizes, so a target symlinked out of the + // models dir has a parent with no draft in it. The sidecar must be + // looked up in the models dir, not beside the canonical file. + // Measured 2026-09-03: serve --speculation dflash ran AR (tau=None) + // on a symlinked qwen3.8-27b.mq5 while the tag form resolved. + let paths = test_paths("dflash-sidecar-symlink"); + fs::create_dir_all(&paths.models).unwrap(); + let elsewhere = paths.root.join("artifacts"); + fs::create_dir_all(&elsewhere).unwrap(); + let real_model = elsewhere.join("qwen3.5-9b.mq4v2.base.hfq"); + fs::write(&real_model, b"model").unwrap(); + std::os::unix::fs::symlink(&real_model, paths.models.join("qwen3.5-9b.mq4")).unwrap(); + let draft_path = paths.models.join("qwen35-9b-dflash-mq4.hfq"); + fs::write(&draft_path, b"draft").unwrap(); + let entry = dflash_sidecar_entry("qwen35-9b-dflash-mq4.hfq"); + let resolved = resolved_with_dflash_mode("on", None); + // What serve/run actually pass: the canonicalized path. + let canonical = fs::canonicalize(paths.models.join("qwen3.5-9b.mq4")).unwrap(); + assert_eq!( + canonical.parent().unwrap(), + elsewhere.canonicalize().unwrap() + ); + let params = load_params( + &resolved, + Some(&entry), + &paths.models, + &canonical, + 64, + Some("q8"), + None, + Some("qwen3.5:9b"), + false, + ) + .unwrap(); + assert_eq!(params["draft"], draft_path.display().to_string()); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + pub(crate) fn load_params_dflash_on_fails_closed_when_sidecar_missing() { + // (c) on + missing file errors with a pull hint naming the tag. + let paths = test_paths("dflash-sidecar-on-missing"); + fs::create_dir_all(&paths.models).unwrap(); + let model_path = paths.models.join("qwen3.5-9b.mq4"); + fs::write(&model_path, b"model").unwrap(); + let entry = dflash_sidecar_entry("qwen35-9b-dflash-mq4.hfq"); + let resolved = resolved_with_dflash_mode("on", None); + let error = load_params( + &resolved, + Some(&entry), + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + Some("qwen3.5:9b"), + false, + ) + .expect_err("on without a pulled draft must fail closed"); + let message = format!("{error:#}"); + assert!(message.contains("qwen35-9b-dflash-mq4.hfq"), "{message}"); + assert!(message.contains("hipfire pull qwen3.5:9b"), "{message}"); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + pub(crate) fn load_params_dflash_auto_runs_ar_when_sidecar_missing() { + // (d) auto + missing file yields no draft and no error. + let paths = test_paths("dflash-sidecar-auto-missing"); + fs::create_dir_all(&paths.models).unwrap(); + let model_path = paths.models.join("qwen3.5-9b.mq4"); + fs::write(&model_path, b"model").unwrap(); + let entry = dflash_sidecar_entry("qwen35-9b-dflash-mq4.hfq"); + let resolved = resolved_with_dflash_mode("auto", None); + let params = load_params( + &resolved, + Some(&entry), + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + Some("qwen3.5:9b"), + false, + ) + .unwrap(); + assert_eq!(params["dflash_mode"], "auto"); + assert!( + params.get("draft").is_none(), + "auto without a pulled draft runs AR" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + pub(crate) fn load_params_explicit_draft_wins_over_dflash_sidecar() { + // (e) developer.dflash_draft beats the sidecar even when pulled. + let paths = test_paths("dflash-sidecar-explicit-wins"); + fs::create_dir_all(&paths.models).unwrap(); + let model_path = paths.models.join("qwen3.5-9b.mq4"); + fs::write(&model_path, b"model").unwrap(); + let draft_path = paths.models.join("qwen35-9b-dflash-mq4.hfq"); + fs::write(&draft_path, b"draft").unwrap(); + let entry = dflash_sidecar_entry("qwen35-9b-dflash-mq4.hfq"); + let explicit = "/tmp/custom-draft.hfq"; + let resolved = resolved_with_dflash_mode("auto", Some(explicit)); + let params = load_params( + &resolved, + Some(&entry), + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + Some("qwen3.5:9b"), + false, + ) + .unwrap(); + assert_eq!(params["draft"], explicit); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + pub(crate) fn load_params_final_off_drops_dflash_sidecar() { + // (f) off never carries the sidecar, and a final off selector drops + // a previously resolved one. + let paths = test_paths("dflash-sidecar-off-drops"); + fs::create_dir_all(&paths.models).unwrap(); + let model_path = paths.models.join("qwen3.5-9b.mq4"); + fs::write(&model_path, b"model").unwrap(); + let draft_path = paths.models.join("qwen35-9b-dflash-mq4.hfq"); + fs::write(&draft_path, b"draft").unwrap(); + let entry = dflash_sidecar_entry("qwen35-9b-dflash-mq4.hfq"); + let resolved = resolved_with_dflash_mode("off", None); + let params = load_params( + &resolved, + Some(&entry), + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + Some("qwen3.5:9b"), + false, + ) + .unwrap(); + assert_eq!(params["dflash_mode"], "off"); + assert!( + params.get("draft").is_none(), + "off must not resolve the sidecar" + ); + + // Resolve under auto, then a final off selector drops it. + let resolved = resolved_with_dflash_mode("auto", None); + let mut params = load_params( + &resolved, + Some(&entry), + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + Some("qwen3.5:9b"), + false, + ) + .unwrap(); + assert_eq!(params["draft"], draft_path.display().to_string()); + apply_speculation_selector(&mut params, "off").unwrap(); + project_dflash_draft(&mut params, developer_dflash_draft(&resolved)); + assert_eq!(params["dflash_mode"], "off"); + assert!( + params.get("draft").is_none(), + "final off must drop the sidecar draft" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + pub(crate) fn load_params_skips_sidecar_for_explicit_cli_draft() { + // `run --model-draft` (projected by the caller after load_params) + // always wins: even `on` must not fail closed on a missing sidecar. + let paths = test_paths("dflash-sidecar-cli-explicit"); + fs::create_dir_all(&paths.models).unwrap(); + let model_path = paths.models.join("qwen3.5-9b.mq4"); + fs::write(&model_path, b"model").unwrap(); + let entry = dflash_sidecar_entry("qwen35-9b-dflash-mq4.hfq"); + let resolved = resolved_with_dflash_mode("on", None); + let params = load_params( + &resolved, + Some(&entry), + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + Some("qwen3.5:9b"), + true, + ) + .unwrap(); + assert!(params.get("draft").is_none()); + fs::remove_dir_all(&paths.root).unwrap(); + } + + /// Minimal in-memory registry for rm tests: (tag, target file, dflash file). + fn rm_test_registry(entries: &[(&str, &str, Option<&str>)]) -> RegistryV1 { + let mut models = BTreeMap::new(); + for (tag, file, dflash) in entries { + models.insert( + (*tag).to_owned(), + ModelEntry { + repo: "test/repo".into(), + file: (*file).to_owned(), + size_gb: 1.0, + min_vram_gb: 1.0, + desc: "rm test".into(), + dflash: dflash.map(|draft| hipfire_registry::Sidecar { + file: draft.into(), + sha256: None, + size_bytes: None, + }), + ..Default::default() + }, + ); + } + RegistryV1 { + schema_version: hipfire_registry::REGISTRY_SCHEMA_VERSION, + generated_at: "test".into(), + _comment: None, + models, + aliases: BTreeMap::new(), + } + } + + #[test] + fn rm_keeps_shared_dflash_sidecar_while_sibling_target_present() { + // The hw-gate regression on PR #686: `qwen3.8:27b`, `qwen3.8:27b-mq4-pro`, + // and `qwen3.8:27b-mq4-xt` all declare `qwen38-27b-dflash-mq4.hfq`. + // Removing one target must keep the sidecar while a sibling declarer's + // target file is still on disk. + let paths = test_paths("rm-shared-sidecar-kept"); + fs::create_dir_all(&paths.models).unwrap(); + for file in [ + "qwen3.8-27b.mq4", + "qwen3.8-27b.mq4-pro", + "qwen38-27b-dflash-mq4.hfq", + ] { + fs::write(paths.models.join(file), b"fixture").unwrap(); + } + let registry = rm_test_registry(&[ + ( + "qwen3.8:27b", + "qwen3.8-27b.mq4", + Some("qwen38-27b-dflash-mq4.hfq"), + ), + ( + "qwen3.8:27b-mq4-pro", + "qwen3.8-27b.mq4-pro", + Some("qwen38-27b-dflash-mq4.hfq"), + ), + ( + "qwen3.8:27b-mq4-xt", + "qwen3.8-27b.mq4-xt", + Some("qwen38-27b-dflash-mq4.hfq"), + ), + ]); + rm_with_registry( + &paths, + ®istry, + RmArgs { + model: "qwen3.8:27b-mq4-pro".into(), + yes: true, + }, + ) + .unwrap(); + assert!( + !paths.models.join("qwen3.8-27b.mq4-pro").exists(), + "removed target is gone" + ); + assert!( + paths.models.join("qwen3.8-27b.mq4").exists(), + "sibling target stays" + ); + assert!( + paths.models.join("qwen38-27b-dflash-mq4.hfq").exists(), + "shared sidecar is kept while a sibling declarer is on disk" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + fn rm_removes_dflash_sidecar_with_last_declaring_target() { + // `qwen3.8:27b` still declares the sidecar in the registry, but its + // target file was never downloaded — a registry row alone must not pin + // the sidecar once the last on-disk declarer is removed. + let paths = test_paths("rm-shared-sidecar-last"); + fs::create_dir_all(&paths.models).unwrap(); + for file in ["qwen3.8-27b.mq4-pro", "qwen38-27b-dflash-mq4.hfq"] { + fs::write(paths.models.join(file), b"fixture").unwrap(); + } + let registry = rm_test_registry(&[ + ( + "qwen3.8:27b", + "qwen3.8-27b.mq4", + Some("qwen38-27b-dflash-mq4.hfq"), + ), + ( + "qwen3.8:27b-mq4-pro", + "qwen3.8-27b.mq4-pro", + Some("qwen38-27b-dflash-mq4.hfq"), + ), + ]); + rm_with_registry( + &paths, + ®istry, + RmArgs { + model: "qwen3.8:27b-mq4-pro".into(), + yes: true, + }, + ) + .unwrap(); + assert!( + !paths.models.join("qwen3.8-27b.mq4-pro").exists(), + "removed target is gone" + ); + assert!( + !paths.models.join("qwen38-27b-dflash-mq4.hfq").exists(), + "sidecar goes with the last on-disk declarer" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + fn rm_without_dflash_declaration_leaves_draft_file_alone() { + // A tag with no dflash declaration keeps master behaviour: its target + // goes, and a draft file it never declared is not an rm target. + let paths = test_paths("rm-no-dflash"); + fs::create_dir_all(&paths.models).unwrap(); + fs::write(paths.models.join("qwen3.8-27b.mq4"), b"fixture").unwrap(); + fs::write(paths.models.join("qwen38-27b-dflash-mq4.hfq"), b"draft").unwrap(); + let registry = rm_test_registry(&[("qwen3.8:27b", "qwen3.8-27b.mq4", None)]); + rm_with_registry( + &paths, + ®istry, + RmArgs { + model: "qwen3.8:27b".into(), + yes: true, + }, + ) + .unwrap(); + assert!( + !paths.models.join("qwen3.8-27b.mq4").exists(), + "removed target is gone" + ); + assert!( + paths.models.join("qwen38-27b-dflash-mq4.hfq").exists(), + "an undeclared draft file is never an rm target" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + fn rm_foreign_same_basename_path_removes_only_that_file() { + // PR #686 hw-gate regression: `hipfire rm /elsewhere/qwen3.6-27b.mq4` + // basename-matched the `qwen3.6:27b` entry and deleted the installed + // target plus its sidecars while the installed target stayed. A path + // that merely shares a basename gets no registry identity: only that + // file goes. + let paths = test_paths("rm-foreign-basename"); + fs::create_dir_all(&paths.models).unwrap(); + for file in ["qwen3.6-27b.mq4", "qwen36-27b-dflash-mq4.hfq"] { + fs::write(paths.models.join(file), b"fixture").unwrap(); + } + let elsewhere = paths.root.join("elsewhere"); + fs::create_dir_all(&elsewhere).unwrap(); + let foreign = elsewhere.join("qwen3.6-27b.mq4"); + fs::write(&foreign, b"lookalike").unwrap(); + let registry = rm_test_registry(&[( + "qwen3.6:27b", + "qwen3.6-27b.mq4", + Some("qwen36-27b-dflash-mq4.hfq"), + )]); + rm_with_registry( + &paths, + ®istry, + RmArgs { + model: foreign.display().to_string(), + yes: true, + }, + ) + .unwrap(); + assert!(!foreign.exists(), "the named foreign file is removed"); + assert!( + paths.models.join("qwen3.6-27b.mq4").exists(), + "installed target stays" + ); + assert!( + paths.models.join("qwen36-27b-dflash-mq4.hfq").exists(), + "installed sidecars stay" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + fn rm_installed_path_removes_target_and_sidecars() { + // The same removal by installed path keeps master behaviour: target + // plus declared sidecars go. + let paths = test_paths("rm-installed-path"); + fs::create_dir_all(&paths.models).unwrap(); + for file in ["qwen3.6-27b.mq4", "qwen36-27b-dflash-mq4.hfq"] { + fs::write(paths.models.join(file), b"fixture").unwrap(); + } + let registry = rm_test_registry(&[( + "qwen3.6:27b", + "qwen3.6-27b.mq4", + Some("qwen36-27b-dflash-mq4.hfq"), + )]); + rm_with_registry( + &paths, + ®istry, + RmArgs { + model: paths.models.join("qwen3.6-27b.mq4").display().to_string(), + yes: true, + }, + ) + .unwrap(); + assert!( + !paths.models.join("qwen3.6-27b.mq4").exists(), + "installed target is gone" + ); + assert!( + !paths.models.join("qwen36-27b-dflash-mq4.hfq").exists(), + "declared sidecar goes with its target" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[cfg(unix)] + #[test] + fn registry_entry_for_path_matches_symlinked_artifact() { + // `qwen3.8-27b.mq4-xt` is a symlink out of the models dir: the input + // and the installed entry file canonicalize to the same target, so + // the entry (and its sidecar) still resolve. A same-basename file + // elsewhere canonicalizes elsewhere and gets no entry. + let paths = test_paths("registry-entry-symlink"); + fs::create_dir_all(&paths.models).unwrap(); + let elsewhere = paths.root.join("qcal"); + fs::create_dir_all(&elsewhere).unwrap(); + let real = elsewhere.join("qwen3.8-27b-weights.mq4"); + fs::write(&real, b"weights").unwrap(); + std::os::unix::fs::symlink(&real, paths.models.join("qwen3.8-27b.mq4-xt")).unwrap(); + let foreign_dir = paths.root.join("foreign"); + fs::create_dir_all(&foreign_dir).unwrap(); + let foreign = foreign_dir.join("qwen3.8-27b.mq4-xt"); + fs::write(&foreign, b"lookalike").unwrap(); + let registry = rm_test_registry(&[("qwen3.8:27b-mq4-xt", "qwen3.8-27b.mq4-xt", None)]); + let installed = paths.models.join("qwen3.8-27b.mq4-xt").display().to_string(); + let (tag, _) = registry_entry_for_path(&paths, ®istry, &installed) + .expect("symlinked installed artifact must match by canonical target"); + assert_eq!(tag, "qwen3.8:27b-mq4-xt"); + assert!( + registry_entry_for_path(&paths, ®istry, &foreign.display().to_string()).is_none(), + "same-basename foreign file gets no registry entry" + ); + assert!( + registry_entry_for_path(&paths, ®istry, "qwen3.8:27b-mq4-xt").is_some(), + "tag form still resolves" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] + fn load_params_dflash_on_fails_closed_for_non_registry_artifact() { + // `on` + a path with no registry entry + no explicit draft fails + // closed with the not-managed message instead of silently running AR + // (Fable's earlier note on unregistered basenames). `auto` on the + // same file still runs AR as a bare artifact. + let paths = test_paths("dflash-on-foreign-path"); + fs::create_dir_all(&paths.models).unwrap(); + let foreign_dir = paths.root.join("elsewhere"); + fs::create_dir_all(&foreign_dir).unwrap(); + let foreign = foreign_dir.join("qwen3.6-27b.mq4"); + fs::write(&foreign, b"lookalike").unwrap(); + assert!( + registry_entry_for_path( + &paths, + &rm_test_registry(&[("qwen3.6:27b", "qwen3.6-27b.mq4", None)]), + &foreign.display().to_string() + ) + .is_none(), + "precondition: foreign path has no entry" + ); + let resolved = resolved_with_dflash_mode("on", None); + let error = load_params( + &resolved, + None, + &paths.models, + &foreign, + 64, + Some("q8"), + None, + None, + false, + ) + .expect_err("on without registry identity must fail closed"); + let message = format!("{error:#}"); + assert!(message.contains("not a registry-managed artifact"), "{message}"); + assert!(message.contains("developer.dflash_draft"), "{message}"); + let resolved = resolved_with_dflash_mode("auto", None); + let params = load_params( + &resolved, + None, + &paths.models, + &foreign, + 64, + Some("q8"), + None, + None, + false, + ) + .unwrap(); + assert_eq!(params["dflash_mode"], "auto"); + assert!( + params.get("draft").is_none(), + "auto on a bare artifact runs AR" + ); + fs::remove_dir_all(&paths.root).unwrap(); + } + + #[test] fn run_spec_dflash_projects_inherited_draft_after_config_off() { // Reviewer case: resolved config leaves DFlash off, but an inherited @@ -6780,7 +7724,18 @@ mod tests { let model_path = PathBuf::from("/tmp/test-model.mq4"); // load_params alone must not carry the draft while config mode is off. - let mut params = load_params(&resolved, None, &model_path, 64, Some("q8"), None).unwrap(); + let mut params = load_params( + &resolved, + None, + &model_path.parent().unwrap(), + &model_path, + 64, + Some("q8"), + None, + None, + false, + ) + .unwrap(); assert_eq!(params["dflash_mode"], "off"); assert!( params.get("draft").is_none(), @@ -9044,6 +9999,121 @@ mod tests { assert_eq!(req.get("max_tokens").and_then(|v| v.as_u64()), Some(128)); } + fn bench_args_for_test(prompt: Vec, prompt_file: Option) -> BenchArgs { + BenchArgs { + model: "qwen:test".to_owned(), + runs: 1, + json: true, + exp: false, + matrix: false, + pp: vec![128], + ctx: vec![128], + tg: 128, + max_tokens: 128, + sustained_tg: None, + sustained_ctx: vec![128], + warmups: 1, + kv_mode: None, + kv_backend: None, + redline: false, + speculation: None, + reasoning_on: false, + concurrency: None, + backend: "both".to_owned(), + workload: "both".to_owned(), + prompt, + prompt_file, + } + } + + #[test] + fn bench_prompt_file_conflicts_with_positional_prompt() { + let err = Cli::try_parse_from([ + "hipfire", + "bench", + "qwen:test", + "--prompt-file", + "prompt.txt", + "hello", + ]) + .unwrap_err(); + assert!( + err.to_string().contains("--prompt-file"), + "conflict error should name the flag: {err}" + ); + } + + #[test] + fn bench_resolve_prompt_keeps_historical_default() { + let args = bench_args_for_test(Vec::new(), None); + assert_eq!( + resolve_bench_prompt(&args).unwrap(), + "Explain the theory of general relativity in simple terms." + ); + } + + #[test] + fn bench_resolve_prompt_joins_positional_words() { + let args = bench_args_for_test(vec!["hello".to_owned(), "world".to_owned()], None); + assert_eq!(resolve_bench_prompt(&args).unwrap(), "hello world"); + } + + #[test] + fn bench_resolve_prompt_file_is_verbatim() { + let path = std::env::temp_dir().join("hipfire-bench-prompt-verbatim.txt"); + // Trailing newline included: the file is read as raw bytes, never trimmed. + std::fs::write(&path, "repeat after me\n").unwrap(); + let args = bench_args_for_test(Vec::new(), Some(path.clone())); + let prompt = resolve_bench_prompt(&args).unwrap(); + std::fs::remove_file(&path).ok(); + assert_eq!(prompt, "repeat after me\n"); + assert_eq!( + bench_prompt_md5(&prompt), + format!("{:x}", md5::compute(b"repeat after me\n")) + ); + } + + #[test] + fn bench_resolve_prompt_rejects_file_and_positional() { + let args = bench_args_for_test(vec!["hello".to_owned()], Some(PathBuf::from("prompt.txt"))); + assert!(resolve_bench_prompt(&args).is_err()); + } + + #[test] + fn bench_prompt_warning_threshold() { + // The default short prompt must warn; 256+ tokens must not. + let short = bench_prompt_warning(24).expect("24 tokens must warn"); + assert!( + short.contains("launch overhead"), + "unexpected text: {short}" + ); + assert!( + short.contains("24"), + "warning should name the count: {short}" + ); + assert!(bench_prompt_warning(255).is_some()); + assert!(bench_prompt_warning(256).is_none()); + assert!(bench_prompt_warning(4400).is_none()); + } + + #[test] + fn bench_prompt_tokens_come_from_prefill_plus_cached() { + // The daemon's done event names the prompt as prefill_tokens (+ any + // prompt-cache hit in cached_tokens); there is no prompt_tokens key. + let done = serde_json::json!({"prefill_tokens": 4400, "cached_tokens": 8}); + assert_eq!(bench_prompt_tokens_from_done(&done), Some(4408)); + let no_cache = serde_json::json!({"prefill_tokens": 24}); + assert_eq!(bench_prompt_tokens_from_done(&no_cache), Some(24)); + let neither = serde_json::json!({"tokens": 128, "prompt_tokens": 99}); + assert_eq!(bench_prompt_tokens_from_done(&neither), None); + } + + #[test] + fn bench_prompt_md5_is_hex_of_prompt_bytes() { + // md5("abc") is a fixed vector; guards against swapping in sha256. + assert_eq!(bench_prompt_md5("abc"), "900150983cd24fb0d6963f7d28e17f72"); + } + #[test] fn http_reasoning_nested_max_tokens_alias_resolves_cap_source() { let resolved = resolve(Vec::::new()).unwrap(); diff --git a/crates/hipfire-cli/src/serve/mod.rs b/crates/hipfire-cli/src/serve/mod.rs index 70d1d8cc6..0253e5f99 100644 --- a/crates/hipfire-cli/src/serve/mod.rs +++ b/crates/hipfire-cli/src/serve/mod.rs @@ -1107,9 +1107,7 @@ impl ServeRuntime { meta: &Mutex, minimum_max_seq: Option, ) -> Result { - let (tag, entry) = self - .registry - .model(model) + let (tag, entry) = crate::registry_entry_for_path(&self.paths, &self.registry, model) .map(|(tag, entry)| (Some(tag.to_owned()), Some(entry))) .unwrap_or((None, None)); let mut path = find_model_path(&self.paths, &self.registry, model); @@ -1143,10 +1141,13 @@ impl ServeRuntime { let mut params = load_params( &resolved, entry, + &self.paths.models, &path, max_tokens, self.kv_override.as_deref(), self.kv_backend_override.as_deref(), + tag.as_deref(), + false, )?; if let Some(tp) = self.tp { params["tp"] = serde_json::json!(tp); @@ -1207,9 +1208,19 @@ impl ServeRuntime { .and_then(serde_json::Value::as_bool) .unwrap_or(false); self.current_max_seq = loaded_max_seq; + // Report the model the way it was requested. A path-form + // request now resolves its registry entry (for sidecars and + // tag policy), but clients — serve_harness's warm probe among + // them — compare `/health.model` against the path they asked + // for; a tag only stands in when the request was a tag. + let served_name = if Path::new(model).is_absolute() || model.contains('/') { + model.to_owned() + } else { + tag.unwrap_or_else(|| model.to_owned()) + }; meta.lock() .unwrap_or_else(|error| error.into_inner()) - .current_model = Some(tag.unwrap_or_else(|| model.to_owned())); + .current_model = Some(served_name); } Ok(resolved) } diff --git a/crates/hipfire-daemon/map.md b/crates/hipfire-daemon/map.md index 77509dd6c..b01eb7ccc 100644 --- a/crates/hipfire-daemon/map.md +++ b/crates/hipfire-daemon/map.md @@ -23,7 +23,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| -| [`src/main.rs`](src/main.rs) | 4,148 | 1 | 0 | +| [`src/main.rs`](src/main.rs) | 4,154 | 1 | 0 | | [`src/slots.rs`](src/slots.rs) | 1,526 | 22 | 14 | ### Public API surface @@ -44,6 +44,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 2 modules · 5,674 lines · 23 public items · 14 tests · 0 examples +- 2 modules · 5,680 lines · 23 public items · 14 tests · 0 examples diff --git a/crates/hipfire-daemon/src/main.rs b/crates/hipfire-daemon/src/main.rs index 23d1dc386..617ba2b2d 100644 --- a/crates/hipfire-daemon/src/main.rs +++ b/crates/hipfire-daemon/src/main.rs @@ -1338,6 +1338,12 @@ fn main() { .and_then(|p| p.get("dspark_conf_threshold")) .and_then(|v| v.as_f64()) .map(|t| t as f32), + // DFlash mirrors mtp: on = fail closed on a missing/unloadable draft. + dflash: match dflash_mode { + "on" => Some(true), + "off" => Some(false), + _ => None, // "auto" → loader default + }, mtp: match mtp_mode.as_str() { "on" => Some(true), "off" => Some(false), diff --git a/crates/hipfire-engine/tests/continuous_batch.rs b/crates/hipfire-engine/tests/continuous_batch.rs index 420bba4ce..eb463ac5f 100644 --- a/crates/hipfire-engine/tests/continuous_batch.rs +++ b/crates/hipfire-engine/tests/continuous_batch.rs @@ -165,13 +165,14 @@ fn batch_eligible_only_qwen_text_single_gpu() { } #[test] -fn batch_eligible_allows_dense_lfm11_and_preserves_qwen() { +fn batch_eligible_refuses_lfm11_and_preserves_qwen() { let _l = begin(); - // LFM dense (arch 11) follows same pure exclusions as Qwen; MoE status is not checked here. - assert!(elig( + // LFM (arch 11) has no servable batch path: the capability is false, so + // even a fully clean request is refused and no batch state is allocated. + assert!(!elig( 11, 1, false, false, false, false, false, false, false, false, true, true, 4 )); - assert!(elig( + assert!(!elig( 11, 1, false, false, false, false, false, false, false, false, true, true, 2 )); // Same pure exclusions as Qwen: B=1, pp!=1, ep, images, tools, stops, spec, adaptive, pflash, history, think. @@ -190,7 +191,7 @@ fn batch_eligible_allows_dense_lfm11_and_preserves_qwen() { assert!(!elig( 11, 1, false, true, false, false, false, false, false, false, true, true, 4 )); - // Unknown arch beside 5/6/11 stays ineligible. + // Archs beside 5/6 stay ineligible (11 included, via the caps refusal above). assert!(!elig( 12, 1, false, false, false, false, false, false, false, false, true, true, 4 )); diff --git a/crates/hipfire-generate/map.md b/crates/hipfire-generate/map.md index f59f344a3..6322df8e0 100644 --- a/crates/hipfire-generate/map.md +++ b/crates/hipfire-generate/map.md @@ -22,12 +22,12 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| -| [`src/ar.rs`](src/ar.rs) | 4,623 | 49 | 0 | +| [`src/ar.rs`](src/ar.rs) | 4,621 | 49 | 0 | | [`src/batch.rs`](src/batch.rs) | 3,465 | 8 | 0 | -| [`src/common.rs`](src/common.rs) | 1,569 | 48 | 1 | -| [`src/dense.rs`](src/dense.rs) | 8,455 | 91 | 4 | +| [`src/common.rs`](src/common.rs) | 1,619 | 49 | 2 | +| [`src/dense.rs`](src/dense.rs) | 8,458 | 91 | 4 | | [`src/lib.rs`](src/lib.rs) | 58 | 7 | 0 | -| [`src/qwen.rs`](src/qwen.rs) | 6,312 | 60 | 1 | +| [`src/qwen.rs`](src/qwen.rs) | 6,403 | 62 | 2 | | [`src/redline.rs`](src/redline.rs) | 4,361 | 49 | 1 | | [`src/vision.rs`](src/vision.rs) | 2,942 | 9 | 8 | @@ -35,10 +35,10 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside - [`src/ar.rs`](src/ar.rs): `arm_fault_after_prefill`, `qwen_ar_route_think_events`, `qwen_ar_route_filter_text`, `QwenArRouteFinish`, `QwenArTerminalCause`, `resolve`, `qwen_ar_finish_route`, `qwen_ar_eos_filter_config`, `qwen_ar_observe_and_route`, `qwen_ar_drain_pending_into_router`, `QwenArRawCommitDisposition`, `qwen_ar_raw_commit_token`, +37 more - [`src/batch.rs`](src/batch.rs): `lfm_prefill_cancellable_or_fallback`, `is_batch_request_eligible`, `drive_qwen_continuous_batch`, `drive_lfm_continuous_batch`, `attach_qwen_ep_batch_receipt_evidence`, `is_qwen_ep_batch_request_eligible`, `drive_qwen35_ep_continuous_batch`, `emit_uncorrelated_error` -- [`src/common.rs`](src/common.rs): `asst_turn_fingerprint`, `strip_think_for_fingerprint`, `normalize_asst_turn_for_fingerprint`, `emit_spec_cancel_after_rollback`, `RollbackEpilogue`, `production_fail_closed_rollback`, `production_fail_closed_rollback_live`, `emit_fail_closed_error`, `ds4_gen_start_contract_version`, `gen_start_contract_version_for_arch`, `Ds4MalformedTerminalAction`, `ds4_malformed_terminal_action`, +36 more +- [`src/common.rs`](src/common.rs): `asst_turn_fingerprint`, `strip_think_for_fingerprint`, `normalize_asst_turn_for_fingerprint`, `emit_spec_cancel_after_rollback`, `RollbackEpilogue`, `production_fail_closed_rollback`, `production_fail_closed_rollback_live`, `emit_fail_closed_error`, `ds4_gen_start_contract_version`, `gen_start_contract_version_for_arch`, `Ds4MalformedTerminalAction`, `ds4_malformed_terminal_action`, +37 more - [`src/dense.rs`](src/dense.rs): `glimmer_turn_key`, `emit_active_attempt_error`, `Ds4SpecWireTerminal`, `ds4_spec_wire_terminal`, `ds4_cache_action`, `ds4_ar_client_abort`, `GlimmerSpecMode`, `glimmer_spec_admission`, `write_error_envelope`, `generate_deepseek4_spec`, `generate_deepseek4`, `ds4_heterogeneous_client_abort`, +79 more - [`src/lib.rs`](src/lib.rs): `common`, `ar`, `qwen`, `dense`, `vision`, `redline`, `batch` -- [`src/qwen.rs`](src/qwen.rs): `EpSampling`, `generate_ep`, `ep_emit_token`, `ep_serve_qwen35_dense_tp`, `ep_emit_done`, `ep_reset_after_abort`, `ep_emit_abort`, `ep_serve_ds4`, `ep_serve_minimax`, `qwen_history_tool_render`, `plan_prompt_cache`, `plan_from_rendered`, +48 more +- [`src/qwen.rs`](src/qwen.rs): `EpSampling`, `EpServeTarget`, `ep_serve_target`, `generate_ep`, `ep_emit_token`, `ep_serve_qwen35_dense_tp`, `ep_emit_done`, `ep_reset_after_abort`, `ep_emit_abort`, `ep_serve_ds4`, `ep_serve_minimax`, `qwen_history_tool_render`, +50 more - [`src/redline.rs`](src/redline.rs): `RedlineQwenSnapshot`, `json`, `RedlineDeepseek4Snapshot`, `RedlineDsparkVerifySnapshot`, `RedlineSnapshot`, `logits`, `kv`, `recurrent`, `gdn_frame`, `redline_qwen_snapshot`, `redline_deepseek4_snapshot`, `RedlineLfm2MoeSnapshot`, +37 more - [`src/vision.rs`](src/vision.rs): `ImageSource`, `GenerateVLParams`, `vl_no_eviction_kv_cap`, `generate_vl`, `generate_vl_dots_ocr`, `decode_vl_dots_ocr_ngram`, `run_dots_ocr_ngram_loop`, `generate_dots_ocr_text`, `generate_lfm2_vl` @@ -55,6 +55,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 8 modules · 31,785 lines · 321 public items · 224 tests · 0 examples +- 8 modules · 31,927 lines · 324 public items · 227 tests · 0 examples diff --git a/crates/hipfire-generate/src/ar.rs b/crates/hipfire-generate/src/ar.rs index f3f9082ee..278e9d5ef 100644 --- a/crates/hipfire-generate/src/ar.rs +++ b/crates/hipfire-generate/src/ar.rs @@ -1160,13 +1160,11 @@ pub fn generate( // eager dense (ModelState::Gemma4) and lowered/MoE // (ModelState::Gemma4Lowered). The generate body is eager-only, so a // lowered load must fail loudly here rather than silently run eager - // against lowered weights. + // against lowered weights. Admission now refuses lowered loads + // before any device allocation; this arm stays as the fail-closed + // net. Message is shared with the admission refusal by construction. if m.gemma4_lowered_mut().is_some() { - emit_error_with_id( - stdout, - id, - "gemma4 lowered/MoE generate not yet wired on this build (eager dense only) — reload without batched/WMMA prefill opt-in or the MoE variant", - ); + emit_error_with_id(stdout, id, hipfire_arch_gemma4::LOWERED_GENERATE_REFUSAL); return; } let _ = ( diff --git a/crates/hipfire-generate/src/common.rs b/crates/hipfire-generate/src/common.rs index 724abff7e..77d49f25b 100644 --- a/crates/hipfire-generate/src/common.rs +++ b/crates/hipfire-generate/src/common.rs @@ -572,6 +572,26 @@ pub fn qwen_dflash_hit_length_cap( generated >= max_tokens && !decoded_eot && !semantic_stop } +/// Shared ctx-capacity margin for the spec entry guard (`generate_dflash` +/// AR fallback) and the in-loop guard (`generate_spec` hard error). A +/// request fits only when prompt + budget + one full draft block fits the +/// draft's context-indexed structures; the `+ block_size` margin is what the +/// mid-loop `position + block_size >= ctx_capacity` break enforces per +/// cycle. Both sites must use this so any request `generate_spec` would +/// refuse falls back to AR at entry instead of erroring after `gen_start` +/// (audit-DFlash Broken 5). +pub fn spec_ctx_request_fits( + prompt_len: usize, + max_tokens: usize, + block_size: usize, + ctx_capacity: usize, +) -> bool { + prompt_len + .saturating_add(max_tokens) + .saturating_add(block_size) + <= ctx_capacity +} + /// Extract held ToolCalls from a FinishSummary (generate_spec holds them). pub fn finish_summary_held_tool_calls( finish: &FinishSummary, @@ -1289,6 +1309,11 @@ pub struct SpecRun { /// `finish.decoded_eot` — wrappers OR both so stop-at-max_tokens wins over /// length. GrammarViolation is fail-closed, not carried here. pub semantic_stop: Option, + /// The mid-loop `position + block_size >= ctx_capacity` break fired: the + /// draft cannot host another full block. Wrappers OR this into the + /// length-cap decision so the turn reports `finish_reason=length` with + /// no cache store instead of a natural `stop`. + pub ctx_exhausted: bool, /// Truthful rollback attestation when this turn ended fail-closed /// (grammar / open-think / malformed). `None` on safe Done paths. pub fail_closed_rollback: Option, @@ -1537,7 +1562,7 @@ pub fn token_logprob_fields( #[cfg(test)] mod tests { - use super::latch_request_think_cap; + use super::{latch_request_think_cap, spec_ctx_request_fits}; #[test] fn numeric_think_cap_latches_once_and_keeps_first_position() { @@ -1566,4 +1591,29 @@ mod tests { assert!(latched); assert_eq!(mark, Some(4096)); } + + #[test] + fn spec_ctx_request_fits_holds_one_block_margin() { + // Sum is prompt + max_tokens + block_size vs ctx_capacity: exactly + // at cap fits, cap+1 refuses, and a bare prompt+max_tokens == cap + // still refuses once the block margin is added (the band the entry + // guard used to admit and the loop guard then rejected). + assert!(spec_ctx_request_fits(100, 900, 24, 1024)); + assert!(spec_ctx_request_fits(100, 899, 24, 1023)); + assert!(!spec_ctx_request_fits(100, 900, 24, 1023)); + assert!(!spec_ctx_request_fits(1000, 24, 24, 1024)); + assert!(!spec_ctx_request_fits(900, 100, 24, 1023)); + // Zero block degrades to the legacy prompt+max_tokens check. + assert!(spec_ctx_request_fits(100, 900, 0, 1000)); + assert!(!spec_ctx_request_fits(100, 901, 0, 1000)); + // Saturating arithmetic: huge budgets clamp instead of panicking + // (debug) or wrapping (release) into a false fit. + assert!(spec_ctx_request_fits(usize::MAX, 1, 1, usize::MAX)); + assert!(!spec_ctx_request_fits( + usize::MAX - 10, + 20, + 0, + usize::MAX - 1 + )); + } } diff --git a/crates/hipfire-generate/src/dense.rs b/crates/hipfire-generate/src/dense.rs index aac549982..6b5aaac55 100644 --- a/crates/hipfire-generate/src/dense.rs +++ b/crates/hipfire-generate/src/dense.rs @@ -559,12 +559,15 @@ pub fn generate_deepseek4_spec( } // Semantic stop (StopSequence/EOS/ThinkCap) or decoded_eot at cap is not // length — preserves stop/tool_calls when generated == max_tokens. - let hit_length_cap = qwen_dflash_hit_length_cap( - run.generated, - max_tokens, - run.finish.decoded_eot, - run.semantic_stop.is_some(), - ); + // A ctx-exhausted mid-loop break is a length stop even when the token + // budget is unspent: same `length` + no-store path as the cap below. + let hit_length_cap = run.ctx_exhausted + || qwen_dflash_hit_length_cap( + run.generated, + max_tokens, + run.finish.decoded_eot, + run.semantic_stop.is_some(), + ); match ds4_spec_wire_terminal( run.finish.finish_reason, run.finish.tool_calls, diff --git a/crates/hipfire-generate/src/qwen.rs b/crates/hipfire-generate/src/qwen.rs index 5bea5b5ba..834900f4a 100644 --- a/crates/hipfire-generate/src/qwen.rs +++ b/crates/hipfire-generate/src/qwen.rs @@ -71,6 +71,29 @@ pub struct EpSampling { pub min_p: Option, } +/// Which EP serve body owns an arch_id. Pure so the dispatch contract is +/// unit-testable. Archs without an `EpArch` (LFM2, Cohere2, anything new) must +/// NOT reach a serve body: the old `_ => ep_serve_ds4` fallthrough ran the +/// DeepSeek4 EP protocol against foreign weights instead of refusing. +#[derive(Debug, Clone, Copy, PartialEq, Eq)] +pub enum EpServeTarget { + Minimax, + Qwen35DenseTp, + Deepseek4, + UnsupportedArch(u32), +} + +/// EP serve dispatch by arch_id. 9/10/5|6 keep their existing servers; +/// anything else is an explicit refusal, never a wrong-server fallthrough. +pub fn ep_serve_target(arch_id: u32) -> EpServeTarget { + match arch_id { + 10 => EpServeTarget::Minimax, + 5 | 6 => EpServeTarget::Qwen35DenseTp, + 9 => EpServeTarget::Deepseek4, + other => EpServeTarget::UnsupportedArch(other), + } +} + pub fn generate_ep( m: &mut LoadedModel, stdout: &mut std::io::Stdout, @@ -225,8 +248,8 @@ pub fn generate_ep( hipfire_loader::EpEosRoute::Qwen35 => m.qwen35_eos_tok, hipfire_loader::EpEosRoute::Deepseek4 => m.deepseek4_eos_tok, }; - match m.arch_id { - 10 => ep_serve_minimax( + match ep_serve_target(m.arch_id) { + EpServeTarget::Minimax => ep_serve_minimax( m, stdout, id, @@ -237,7 +260,7 @@ pub fn generate_ep( primed_think, sampling, ), - 5 | 6 => ep_serve_qwen35_dense_tp( + EpServeTarget::Qwen35DenseTp => ep_serve_qwen35_dense_tp( m, stdout, id, @@ -249,7 +272,7 @@ pub fn generate_ep( primed_think, sampling, ), - _ => ep_serve_ds4( + EpServeTarget::Deepseek4 => ep_serve_ds4( m, stdout, id, @@ -261,6 +284,26 @@ pub fn generate_ep( stop, sampling, ), + EpServeTarget::UnsupportedArch(arch) => { + // Fail closed: no EpArch exists for this arch_id (LFM2, Cohere2, + // anything new), so there is no correct server to call. Name the + // arch rather than running another arch's EP protocol against + // foreign weights. + emit_active_attempt_error( + stdout, + Some(id), + &format!( + "EP generate not supported for arch_id={arch} \ + (only 9/DeepSeek4, 10/MiniMax and dense 5|6 Qwen3.5 \ + have an EP serve path)" + ), + "unsupported", + false, + false, + ); + let _ = stdout.flush(); + return; + } } } @@ -2075,7 +2118,16 @@ pub fn generate_dflash( .as_ref() .map(|s| s.ctx_capacity()) .unwrap_or(usize::MAX); - if prompt_tokens.len().saturating_add(max_tokens) > spec_ctx_capacity { + let spec_block_size = m.speculator.as_ref().map(|s| s.block_size()).unwrap_or(0); + // Shared margin with the `generate_spec` hard guard below: prompt + + // budget + one draft block must fit, so any request the loop would refuse + // falls back to AR here instead of erroring after `gen_start`. + if !spec_ctx_request_fits( + prompt_tokens.len(), + max_tokens, + spec_block_size, + spec_ctx_capacity, + ) { emit_qwen_ar_info( stdout, id, @@ -2430,8 +2482,10 @@ pub fn generate_dflash( im_end_token, ); let semantic_stop = run.semantic_stop.is_some(); - let hit_length_cap = - qwen_dflash_hit_length_cap(run.generated, max_tokens, decoded_eot, semantic_stop); + // A ctx-exhausted mid-loop break is a length stop even when the token + // budget is unspent: same `length` + no-store path as the cap below. + let hit_length_cap = run.ctx_exhausted + || qwen_dflash_hit_length_cap(run.generated, max_tokens, decoded_eot, semantic_stop); // Prefer producer-visible channel; fall back to finish Token events. let visible = if !run.finish.visible_text.is_empty() { run.finish.visible_text.clone() @@ -2637,12 +2691,13 @@ pub fn generate_dflash( let emit_tool_calls = extract_tool_calls_from_text(&decoded_full); // Semantic stop / decoded_eot at the budget boundary is stop/tool_calls, // not length — same rule as the qwen_semantic_v2 path. - let hit_length_cap = qwen_dflash_hit_length_cap( - run.generated, - max_tokens, - run.finish.decoded_eot, - run.semantic_stop.is_some(), - ); + let hit_length_cap = run.ctx_exhausted + || qwen_dflash_hit_length_cap( + run.generated, + max_tokens, + run.finish.decoded_eot, + run.semantic_stop.is_some(), + ); let finish_reason = if hit_length_cap { "length" } else if !emit_tool_calls.is_empty() { @@ -2977,12 +3032,11 @@ pub fn generate_spec( let _ = stdout.flush(); return None; } + // Shared margin with the `generate_dflash` entry fallback above (same + // predicate): without eviction the entry already diverted these to AR, + // so this is belt-and-suspenders for direct `generate_spec` callers. if m.eviction.is_none() - && prompt_tokens - .len() - .saturating_add(max_tokens) - .saturating_add(block_size) - > ctx_capacity + && !spec_ctx_request_fits(prompt_tokens.len(), max_tokens, block_size, ctx_capacity) { emit_active_attempt_error( stdout, @@ -3110,15 +3164,25 @@ pub fn generate_spec( let mut emit: Box = match carrier.make_spec_emitter(emit_ctx) { Ok(e) => e, Err(e) => { - emit_active_attempt_error( - stdout, - Some(id), - &format!("{}", e), - "validation", - false, - false, + // Post-prefill failure: the target's KV/DeltaNet/drafter hidden + // already advanced and host seq_pos/conversation_tokens were + // cleared on a cold start. Fail closed like every other + // post-prefill error exit (prefill/step Err, realign, forced + // terminal): live rollback first, then one correlated error — + // otherwise the next turn LCPs against a dirty GPU + // (audit-Dflash Broken 4). + let msg = format!("make_spec_emitter: {e}"); + let ep = production_fail_closed_rollback_live( + &mut m.seq_pos, + &mut m.conversation_tokens, + &mut m.prefill_checkpoints, + &mut m.dflash_checkpoints, + &mut m.asst_turn_cache, + gpu, + slot, + spec.as_mut(), ); - let _ = stdout.flush(); + emit_fail_closed_error(stdout, Some(id), &msg, "validation", false, &ep); return None; } }; @@ -3141,6 +3205,11 @@ pub fn generate_spec( let mut spec_cycles = 0usize; let mut spec_accepted = 0usize; let mut generated = 0usize; + // Set when the mid-loop `position + block_size >= ctx_capacity` break + // fires: the draft's context-indexed structures cannot host another + // full block, so the epilogue must report a length stop + // (`finish_reason=length`, no cache store) rather than a natural stop. + let mut ctx_exhausted = false; // Post-prefill compaction (FlashCASK pattern from dflash_spec_demo). // If the prompt already filled past budget+beta, compact once before @@ -3356,6 +3425,7 @@ pub fn generate_spec( return None; } if position.saturating_add(block_size) >= ctx_capacity { + ctx_exhausted = true; break; } @@ -3948,6 +4018,7 @@ pub fn generate_spec( finish, grammar_violated, semantic_stop, + ctx_exhausted, fail_closed_rollback, prefill_s: t_prefill.duration_since(t0).as_secs_f64(), total_s: t_end.duration_since(t0).as_secs_f64(), @@ -6310,3 +6381,23 @@ mod deepseek4_reasoning_prefix_tests { } // --- iter appended --- + +#[cfg(test)] +mod ep_serve_target_tests { + use super::{ep_serve_target, EpServeTarget}; + + #[test] + fn known_ep_archs_keep_their_servers_but_all_others_refuse() { + // Adjacent supported: DS4, MiniMax and dense Qwen3.5 keep their servers. + assert_eq!(ep_serve_target(9), EpServeTarget::Deepseek4); + assert_eq!(ep_serve_target(10), EpServeTarget::Minimax); + assert_eq!(ep_serve_target(5), EpServeTarget::Qwen35DenseTp); + assert_eq!(ep_serve_target(6), EpServeTarget::Qwen35DenseTp); + // No EpArch exists for these: explicit refusal naming the arch, + // never the old wrong-server DS4 fallthrough. + assert_eq!(ep_serve_target(11), EpServeTarget::UnsupportedArch(11)); + assert_eq!(ep_serve_target(12), EpServeTarget::UnsupportedArch(12)); + assert_eq!(ep_serve_target(13), EpServeTarget::UnsupportedArch(13)); + assert_eq!(ep_serve_target(0), EpServeTarget::UnsupportedArch(0)); + } +} diff --git a/crates/hipfire-generate/tests/continuous_batch.rs b/crates/hipfire-generate/tests/continuous_batch.rs index cfc4d604d..fa31dee68 100644 --- a/crates/hipfire-generate/tests/continuous_batch.rs +++ b/crates/hipfire-generate/tests/continuous_batch.rs @@ -156,13 +156,14 @@ fn batch_eligible_only_qwen_text_single_gpu() { } #[test] -fn batch_eligible_allows_dense_lfm11_and_preserves_qwen() { +fn batch_eligible_refuses_lfm11_and_preserves_qwen() { let _l = begin(); - // LFM dense (arch 11) follows same pure exclusions as Qwen; MoE status is not checked here. - assert!(elig( + // LFM (arch 11) has no servable batch path: the capability is false, so + // even a fully clean request is refused and no batch state is allocated. + assert!(!elig( 11, 1, false, false, false, false, false, false, false, false, true, true, 4 )); - assert!(elig( + assert!(!elig( 11, 1, false, false, false, false, false, false, false, false, true, true, 2 )); // Same pure exclusions as Qwen: B=1, pp!=1, ep, images, tools, stops, spec, adaptive, pflash, history, think. @@ -181,7 +182,7 @@ fn batch_eligible_allows_dense_lfm11_and_preserves_qwen() { assert!(!elig( 11, 1, false, true, false, false, false, false, false, false, true, true, 4 )); - // Unknown arch beside 5/6/11 stays ineligible. + // Archs beside 5/6 stay ineligible (11 included, via the caps refusal above). assert!(!elig( 12, 1, false, false, false, false, false, false, false, false, true, true, 4 )); diff --git a/crates/hipfire-generate/tests/qwen_dflash_ctx_exhausted_tests.rs b/crates/hipfire-generate/tests/qwen_dflash_ctx_exhausted_tests.rs new file mode 100644 index 000000000..d94acb29c --- /dev/null +++ b/crates/hipfire-generate/tests/qwen_dflash_ctx_exhausted_tests.rs @@ -0,0 +1,71 @@ +// SPDX-License-Identifier: Apache-2.0 +// Copyright (c) 2026 Kaden Schutt +// hipfire — see LICENSE and NOTICE in the project root. + +//! Draft-context exhaustion terminal contract. +//! +//! A mid-loop `position + block_size >= ctx_capacity` break with +//! `generated < max_tokens` must report `length` (no tool release, no cache +//! store), not a natural `stop`. Kept separate from +//! `qwen_dflash_semantic_terminal_tests.rs` so that file's rustfmt debt is not +//! dragged into this change. + +use hipfire_generate::common::qwen_dflash_hit_length_cap; +use hipfire_generate::qwen::{ + qwen_dflash_apply_cache_action, qwen_dflash_cache_action, qwen_dflash_wire_terminal, + QwenDflashWireTerminal, +}; +use hipfire_runtime::prompt_frame::ToolCall; +use hipfire_runtime::spec::{ClientEvent, FinishSummary}; + +fn summary_tool_calls(calls: Vec) -> FinishSummary { + let n = calls.len(); + FinishSummary { + events: vec![ClientEvent::ToolCalls(calls)], + finish_reason: "tool_calls", + tool_calls: n, + visible_text: "Sure.".into(), + decoded_eot: false, + open_think: false, + } +} + +#[test] +fn ctx_exhausted_maps_to_length_with_budget_unspent() { + let calls = vec![ToolCall { + id: None, + name: "t".into(), + arguments: serde_json::json!({}), + rendered_body: None, + }]; + let fin = summary_tool_calls(calls); + // The budget alone would not have stopped this turn. + assert!(!qwen_dflash_hit_length_cap(10, 16, false, false)); + // Wrapper-level mapping (`generate_dflash` / dense spec epilogue): the + // `ctx_exhausted` flag is OR-ed into the length decision. + let ctx_exhausted = true; + let hit_length_cap = ctx_exhausted || qwen_dflash_hit_length_cap(10, 16, false, false); + let term = qwen_dflash_wire_terminal(&fin, hit_length_cap, false, "partial", false); + match &term { + QwenDflashWireTerminal::Done { + finish_reason, + release_tool_calls, + store_cache, + wire_tool_calls, + fingerprint_text, + } => { + assert_eq!(*finish_reason, "length"); + assert!(!*release_tool_calls); + assert!(!*store_cache); + assert!(wire_tool_calls.is_empty()); + assert!(fingerprint_text.is_empty()); + } + other => panic!("expected length Done, got {other:?}"), + } + let action = qwen_dflash_cache_action(&term); + assert!(!action.store); + assert!( + qwen_dflash_apply_cache_action(|_, _| panic!("must not insert"), &action, vec![1, 2]) + .is_none() + ); +} diff --git a/crates/hipfire-loader/map.md b/crates/hipfire-loader/map.md index b563b6c87..142a476aa 100644 --- a/crates/hipfire-loader/map.md +++ b/crates/hipfire-loader/map.md @@ -24,15 +24,15 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| | [`src/batch_staging.rs`](src/batch_staging.rs) | 336 | 4 | 0 | -| [`src/carriers.rs`](src/carriers.rs) | 2,516 | 11 | 3 | -| [`src/lib.rs`](src/lib.rs) | 4,838 | 95 | 22 | +| [`src/carriers.rs`](src/carriers.rs) | 2,521 | 11 | 3 | +| [`src/lib.rs`](src/lib.rs) | 4,978 | 97 | 25 | | [`src/spec_build.rs`](src/spec_build.rs) | 233 | 4 | 0 | ### Public API surface - [`src/batch_staging.rs`](src/batch_staging.rs): `BatchStaging`, `qwen_batch_weight_formats_supported`, `qwen_ep_batch_weight_formats_supported`, `stage_continuous_batch` - [`src/carriers.rs`](src/carriers.rs): `Qwen2Carrier`, `Qwen35Carrier`, `LlamaCarrier`, `DotsOcrCarrier`, `Deepseek4Carrier`, `MinimaxCarrier`, `Lfm2MoeCarrier`, `Cohere2MoeCarrier`, `MapleCarrier`, `Gemma4Carrier`, `MuseGlimmerCarrier` -- [`src/lib.rs`](src/lib.rs): `batch_staging`, `carriers`, `spec_build`, `Carrier`, `carrier_for`, `ContinuousBatchRoute`, `continuous_batch_route`, `BenchDecodeRoute`, `bench_decode_route`, `VisionRoute`, `vision_route`, `EpPromptRoute`, +83 more +- [`src/lib.rs`](src/lib.rs): `batch_staging`, `carriers`, `spec_build`, `Carrier`, `carrier_for`, `ContinuousBatchRoute`, `continuous_batch_route`, `BenchDecodeRoute`, `bench_decode_route`, `VisionRoute`, `vision_route`, `EpPromptRoute`, +85 more - [`src/spec_build.rs`](src/spec_build.rs): `Qwen35SlotGuard`, `take`, `model_slot`, `build_speculator` ### Dependencies (from `Cargo.toml`) @@ -48,6 +48,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 4 modules · 7,923 lines · 114 public items · 25 tests · 1 examples +- 4 modules · 8,068 lines · 116 public items · 28 tests · 1 examples diff --git a/crates/hipfire-loader/src/carriers.rs b/crates/hipfire-loader/src/carriers.rs index c7dba2dba..b7e2fa913 100644 --- a/crates/hipfire-loader/src/carriers.rs +++ b/crates/hipfire-loader/src/carriers.rs @@ -1467,7 +1467,12 @@ impl Carrier for Lfm2MoeCarrier { } fn caps(&self) -> saddle_core::caps::ArchCaps { saddle_core::caps::ArchCaps { - supports_continuous_batch: true, + // Continuous batching is NOT servable: the generate-side eligibility + // (`is_batch_request_eligible`) returns false unconditionally for + // LFM, so staging a batch state only spends VRAM on state that is + // never driven. Declare false so the route never admits it and the + // state is never allocated. Single-stream LFM is unaffected. + supports_continuous_batch: false, supports_ep_batch: false, dflash: None, supports_mtp: false, diff --git a/crates/hipfire-loader/src/lib.rs b/crates/hipfire-loader/src/lib.rs index c099ef281..f20fd1535 100644 --- a/crates/hipfire-loader/src/lib.rs +++ b/crates/hipfire-loader/src/lib.rs @@ -188,12 +188,14 @@ pub enum ContinuousBatchRoute { Qwen35, Lfm2Moe, } -/// Exact arch_id -> continuous-batch route. Mirrors the two batch-capable -/// families (qwen35 5|6, lfm2moe 11). No carrier probing — pure id match. +/// Exact arch_id -> continuous-batch route. Only qwen35 5|6 admits: LFM2 (11) +/// has no servable batch path (see the lfm2moe carrier caps), so the route +/// refuses it and no batch state is ever allocated. No carrier probing — pure +/// id match. `ContinuousBatchRoute::Lfm2Moe` stays for the staging body, which +/// is now unreachable. pub fn continuous_batch_route(arch_id: u32) -> Option { match arch_id { 5 | 6 => Some(ContinuousBatchRoute::Qwen35), - 11 => Some(ContinuousBatchRoute::Lfm2Moe), _ => None, } } @@ -1676,6 +1678,14 @@ pub(crate) fn parse_state_quant( // ─── Core arch carrier load ───────────────────────────────────────────── /// Hard-error free for unfinished qwen35 finish path: bundle + optional VL. +/// +/// `free_qwen35_bundle` returns every buffer to the GPU pool; only a drain +/// hands the VRAM back to the driver, and `unload_model` is normally the +/// one that drains. A load that fails here never reaches `unload_model`, +/// so without the drain the whole target (~15 GB on a 27B) stayed pooled — +/// the hw-gate Fable seat measured ~5 GB retained after a refused +/// `dflash_mode=on` load on top of the next resident model, compounding on +/// every lazy serve retry. Mirror `unload_model`: invalidate graphs, drain. fn rollback_unfinished_qwen35( err: String, bundle: Qwen35Bundle, @@ -1689,6 +1699,8 @@ fn rollback_unfinished_qwen35( if let Some(vw) = vision_weights { vw.free_gpu(gpu); } + gpu.invalidate_graph_state(); + gpu.drain_pool(); if notes.is_empty() { err } else { @@ -1965,6 +1977,16 @@ fn finish_qwen35_load( Some(s) } Err(e) => { + if ctx.spec.dflash == Some(true) { + return Err(rollback_unfinished_qwen35( + format!( + "DFlash draft required (dflash_mode=on) but failed to load ({dp}): {e}" + ), + bundle, + vision_weights, + ctx.gpu, + )); + } eprintln!( " DFlash draft load failed ({}): {} — falling back to AR only", dp, e @@ -2841,8 +2863,46 @@ impl Drop for Qwen35DenseTpStaging { } } +/// Admission refusal for Qwen3.5-MoE under expert-parallel load (#683 family). +/// Pure so the contract is unit-testable: any Qwen3.5 config with routed +/// experts (`num_experts > 0`, i.e. arch 6 and any mis-stamped arch 5) has no +/// EP serve path — `generate_ep` routes arch 6 at the dense-TP server, which +/// only accepts `EpArch::Qwen35DenseTp`. Refuse here, before `Gpus::init_tp` +/// and the per-rank weight upload, instead of after a full 4-rank load. +/// Dense Qwen3.5 (`num_experts == 0`) is unaffected and keeps its EP path. +pub fn qwen35_ep_moe_refusal(arch_id: u32, num_experts: usize) -> Option { + if num_experts > 0 { + Some(format!( + "Qwen3.5-MoE (arch_id={arch_id}) has no EP serve path; use TP or single-GPU" + )) + } else { + None + } +} + +/// Message constructor shared by [`ep_admission`] and the per-entry match +/// backstops so the refusal text cannot drift between the two. +fn ep_unsupported_arch_message(arch_id: u32) -> String { + format!( + "EP not supported for arch_id={arch_id} (expected 5|6 for Qwen3.5, 9 for DeepSeek V4 or 10 for MiniMax)" + ) +} + +/// EP load admission by arch_id. Only archs with an `EpArch` variant may enter +/// expert-parallel load (9/DeepSeek4, 10/MiniMax, 5|6/Qwen3.5); LFM2 (11), +/// Cohere2 (12) and anything else must fail here — right after the host-side +/// HFQ probe, before any device init — otherwise they would reach +/// `generate_ep` with no correct server. Pure so the contract is +/// unit-testable; both `load_model_ep_*` entries call it before dispatching. +pub fn ep_admission(arch_id: u32) -> Result<(), String> { + match arch_id { + 5 | 6 | 9 | 10 => Ok(()), + id => Err(ep_unsupported_arch_message(id)), + } +} + /// Expert-parallel (EP) model load — shards the routed experts across `tp` ranks -/// (`Gpus::init_tp` + per-arch sharded weight load), wrapped in a staging guard so +/// (`Gpus::init_ep` + per-arch sharded weight load), wrapped in a staging guard so /// a mid-load failure frees every already-loaded rank's VRAM (no leak, prior model /// at the call site left intact). ds4 (arch_id 9) and MiniMax (arch_id 10) only. /// @@ -2909,6 +2969,8 @@ pub fn load_model_ep_with_kv_mode( state_quant: Option<&str>, ) -> Result { let hfq = HfqFile::open(Path::new(path)).map_err(|e| format!("{e}"))?; + // Admission: refuse archs with no `EpArch` before any per-arch device init. + ep_admission(hfq.arch_id)?; let kv_backend_raw = kv_backend.unwrap_or("contiguous"); let kv_backend_kind: KvBackend = kv_backend_raw.parse().map_err(|err| format!("{err}"))?; match hfq.arch_id { @@ -2926,9 +2988,10 @@ pub fn load_model_ep_with_kv_mode( Err(format!("KV backend '{kv_backend_raw}' requires tp=1")) } 5 | 6 => load_model_ep_qwen35(path, max_seq, tp, kv_mode, kv_backend, state_quant), - id => Err(format!( - "EP not supported for arch_id={id} (expected 5|6 for Qwen3.5, 9 for DeepSeek V4 or 10 for MiniMax)" - )), + // Backstop: `ep_admission` above already refused these; route through the + // shared constructor (not `unreachable!`) so the refusal survives a + // future edit that drops the early call. + id => Err(ep_unsupported_arch_message(id)), } } @@ -2942,6 +3005,8 @@ pub fn load_model_ep_with_compressor_cache( compressor_cache: hipfire_config::Deepseek4CompressorCache, ) -> Result { let hfq = HfqFile::open(Path::new(path)).map_err(|e| format!("{e}"))?; + // Admission: refuse archs with no `EpArch` before any per-arch device init. + ep_admission(hfq.arch_id)?; match hfq.arch_id { 9 => load_model_ep_ds4(path, max_seq, tp, compressor_cache), 10 if compressor_cache == hipfire_config::Deepseek4CompressorCache::F32 => { @@ -2951,10 +3016,13 @@ pub fn load_model_ep_with_compressor_cache( 5 | 6 if compressor_cache == hipfire_config::Deepseek4CompressorCache::F32 => { load_model_ep_qwen35(path, max_seq, tp, None, None, None) } - 5 | 6 => Err("DeepSeek V4 compressor-cache storage cannot be applied to Qwen3.5".to_string()), - id => Err(format!( - "EP not supported for arch_id={id} (expected 5|6 for Qwen3.5, 9 for DeepSeek V4 or 10 for MiniMax)" - )), + 5 | 6 => { + Err("DeepSeek V4 compressor-cache storage cannot be applied to Qwen3.5".to_string()) + } + // Backstop: `ep_admission` above already refused these; route through the + // shared constructor (not `unreachable!`) so the refusal survives a + // future edit that drops the early call. + id => Err(ep_unsupported_arch_message(id)), } } @@ -2991,11 +3059,11 @@ fn load_model_ep_ds4( let rec = hfq.recommended_sampling(); let gpus = - Gpus::init_tp(tp, config.num_hidden_layers).map_err(|e| format!("init_tp: {e:?}"))?; + Gpus::init_ep(tp, config.num_hidden_layers).map_err(|e| format!("init_ep: {e:?}"))?; let n = gpus.devices.len(); if n != tp { return Err(format!( - "init_tp gave {n} devices, expected tp={tp} (check ROCR_VISIBLE_DEVICES / HIP_VISIBLE_DEVICES)" + "init_ep gave {n} devices, expected tp={tp} (check ROCR_VISIBLE_DEVICES / HIP_VISIBLE_DEVICES)" )); } eprintln!("[loader] EP load: tp={tp} arch=ds4 experts={n_exp} (rank r owns e%{tp}==r)"); @@ -3219,11 +3287,11 @@ fn load_model_ep_minimax(path: &str, max_seq: usize, tp: usize) -> Result Result<( } } +#[cfg(test)] +mod ep_admission_tests { + use super::{ep_admission, qwen35_ep_moe_refusal}; + + #[test] + fn qwen35_moe_ep_refuses_before_load_but_dense_admits() { + // Arch 6 MoE under EP: refused with the combination named. + let err = + qwen35_ep_moe_refusal(6, 128).expect("arch-6 MoE + EP must be refused at admission"); + assert!(err.contains("Qwen3.5-MoE"), "reason names the model: {err}"); + assert!(err.contains("no EP serve path"), "reason: {err}"); + assert!(err.contains('6'), "reason names the arch: {err}"); + // Mis-stamped arch 5 with routed experts: same missing serve path. + assert!(qwen35_ep_moe_refusal(5, 128).is_some()); + // Adjacent supported: dense Qwen3.5 (either arch id) keeps its EP path. + assert_eq!(qwen35_ep_moe_refusal(5, 0), None); + assert_eq!(qwen35_ep_moe_refusal(6, 0), None); + } + + #[test] + fn ep_without_eparch_refuses_but_served_archs_admit() { + // LFM2 (11), Cohere2 (12) and anything else with no `EpArch` variant: + for arch in [11u32, 12, 13, 0, 99] { + let err = match ep_admission(arch) { + Ok(()) => panic!("arch {arch} + EP must refuse"), + Err(e) => e, + }; + assert!(err.contains("EP not supported"), "reason: {err}"); + assert!( + err.contains(&arch.to_string()), + "reason names the arch: {err}" + ); + } + // Adjacent supported: DS4, MiniMax and Qwen3.5 keep their EP entries. + for arch in [5u32, 6, 9, 10] { + assert!(ep_admission(arch).is_ok(), "arch {arch} + EP must admit"); + } + } +} + +#[cfg(test)] +mod lfm2_batch_admission_tests { + #[test] + fn lfm2_continuous_batch_never_admits_but_qwen_still_does() { + use super::{carrier_for, continuous_batch_route}; + // Arch 11: the route refuses, so staging takes the fallback arm and + // no Lfm2DecodeBatchState is ever allocated; the caps gate in + // `is_batch_request_eligible` (and the engine scheduler) agrees. + assert_eq!(continuous_batch_route(11), None); + let caps = carrier_for(11).expect("lfm2moe carrier").caps(); + assert!( + !caps.supports_continuous_batch, + "lfm2moe caps must stay false while no batch path is servable" + ); + // Adjacent supported: qwen35 5|6 still admit continuous batching. + assert!(continuous_batch_route(5).is_some()); + assert!(continuous_batch_route(6).is_some()); + assert!( + carrier_for(5) + .expect("qwen35 carrier") + .caps() + .supports_continuous_batch + ); + } +} + #[cfg(test)] mod registry_tests { use super::{resolve_deepseek4_compressor_cache_kv_mode, REGISTRY}; @@ -4164,10 +4303,12 @@ mod registry_tests { } ); assert_eq!(caps_of("minimax"), text_only); + // lfm2moe declares supports_continuous_batch: false — the batch state + // was allocated and never driven (eligibility always false), so the + // capability is truthful only when false. Single-stream LFM unaffected. assert_eq!( caps_of("lfm2moe"), ArchCaps { - supports_continuous_batch: true, supports_images: true, ..text_only } @@ -4189,13 +4330,12 @@ mod registry_tests { } ); - // ── continuous_batch_route: 5|6 -> Qwen35, 11 -> Lfm2Moe ── + // ── continuous_batch_route: 5|6 -> Qwen35 only (11/LFM2 refuses) ── // The Some/None half duplicates caps().supports_continuous_batch; the // variant picks between two distinct staging bodies in batch_staging. for id in 0u32..=14 { let want = match id { 5 | 6 => Some(ContinuousBatchRoute::Qwen35), - 11 => Some(ContinuousBatchRoute::Lfm2Moe), _ => None, }; assert_eq!( diff --git a/crates/hipfire-registry/map.md b/crates/hipfire-registry/map.md index c101c3dfc..13dcf2581 100644 --- a/crates/hipfire-registry/map.md +++ b/crates/hipfire-registry/map.md @@ -22,7 +22,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| -| [`src/lib.rs`](src/lib.rs) | 1,796 | 25 | 18 | +| [`src/lib.rs`](src/lib.rs) | 1,843 | 25 | 19 | ### Public API surface @@ -41,6 +41,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 1 modules · 1,796 lines · 25 public items · 18 tests · 0 examples +- 1 modules · 1,843 lines · 25 public items · 19 tests · 0 examples diff --git a/crates/hipfire-registry/src/lib.rs b/crates/hipfire-registry/src/lib.rs index bed694249..fbce30af5 100644 --- a/crates/hipfire-registry/src/lib.rs +++ b/crates/hipfire-registry/src/lib.rs @@ -188,6 +188,8 @@ pub struct ModelEntry { #[serde(default)] pub dspark: Option, #[serde(default)] + pub dflash: Option, + #[serde(default)] pub default_tool_format: Option, #[serde(default)] pub default_kv_mode: Option, @@ -364,7 +366,7 @@ impl RegistryV1 { return Err(fail(format!("model '{tag}' has invalid size metadata"))); } validate_digest(entry.sha256.as_deref(), tag).map_err(fail)?; - for sidecar in [&entry.triattn, &entry.mtp, &entry.dspark] + for sidecar in [&entry.triattn, &entry.mtp, &entry.dspark, &entry.dflash] .into_iter() .flatten() { @@ -422,6 +424,12 @@ impl RegistryV1 { if self.models.contains_key(&qwen) { return qwen; } + // A bare file name matches its entry. A path never does: matching on + // `file_name()` let a lookalike file outside the models directory + // inherit the installed artifact's identity (sidecars, kv/max_seq + // policy, rm targets). Callers that know the models directory resolve + // paths with canonical comparison instead (`registry_entry_for_path` + // in hipfire-cli); a bare file name carries no directory to confuse. self.models .iter() .find_map(|(tag, entry)| { @@ -430,6 +438,17 @@ impl RegistryV1 { .unwrap_or(normalized) } + /// Exact `entry.file` match for callers that already established the input + /// is the installed artifact (e.g. via canonical path comparison against + /// the models directory). Unlike [`RegistryV1::model`], this never applies + /// tag/alias normalization: `file_name` must be the bare file as stored. + pub fn entry_for_file(&self, file_name: &str) -> Option<(&str, &ModelEntry)> { + self.models + .iter() + .find(|(_, entry)| entry.file == file_name) + .map(|(tag, entry)| (tag.as_str(), entry)) + } + pub fn model(&self, input: &str) -> Option<(&str, &ModelEntry)> { let tag = self.resolve_tag(input); self.models @@ -1028,6 +1047,33 @@ mod tests { "effort-native: absence means uncapped" ); } + #[test] + fn bundled_dflash_sidecars_name_pullable_files() { + // Every `dflash.file` must name a file that some registry entry's + // `file` also names, so `hipfire pull ` fetching the sidecar + // always lands a file that `hipfire pull -draft` could fetch too. + let registry = bundled().unwrap(); + let files: std::collections::BTreeSet<&str> = registry + .models + .values() + .map(|entry| entry.file.as_str()) + .collect(); + let mut paired = 0; + for (tag, entry) in ®istry.models { + if let Some(sidecar) = entry.dflash.as_ref() { + assert!( + files.contains(sidecar.file.as_str()), + "model '{tag}' declares dflash sidecar '{}' with no matching entry file", + sidecar.file + ); + paired += 1; + } + } + assert!( + paired > 0, + "bundled registry should pair at least one dflash sidecar" + ); + } #[test] fn aliases_and_filenames_resolve_to_canonical_tags() { @@ -1043,6 +1089,21 @@ mod tests { "qwen3.8:27b-mq4-xt" ); assert_eq!(registry.resolve_tag("qwen3.8:fast"), "qwen3.8:27b-mq4-xt"); + // A path is not a tag: even a path into the models directory resolves + // to itself here. Callers that know the models directory establish + // identity with canonical path comparison (`registry_entry_for_path` + // in hipfire-cli), so `serve --model ~/.hipfire/models/` still + // sees the entry's sidecars without letting a same-basename lookalike + // elsewhere inherit them (PR #686 hw-gate regression). + assert_eq!( + registry.resolve_tag("/home/u/.hipfire/models/qwen3.8-27b.mq5"), + "/home/u/.hipfire/models/qwen3.8-27b.mq5" + ); + assert_eq!( + registry.resolve_tag("/home/u/.hipfire/models/qwen3.8-27b.mq4"), + "/home/u/.hipfire/models/qwen3.8-27b.mq4" + ); + assert!(registry.model("/home/u/.hipfire/models/qwen3.8-27b.mq4").is_none()); assert_eq!(registry.resolve_tag("deepseek4"), "deepseek-v4-flash"); assert_eq!(registry.resolve_tag("deepseek4:0731"), "deepseek-v4-flash"); @@ -1101,6 +1162,48 @@ mod tests { ); } + #[test] + fn path_inputs_never_resolve_to_registry_tags() { + // PR #686 hw-gate regression: `hipfire rm /elsewhere/qwen3.6-27b.mq4` + // basename-matched the `qwen3.6:27b` entry and deleted the installed + // sidecars. A path — anywhere, including the models directory itself — + // is not a tag and resolves to itself. + let raw = r#"{ + "schema_version":1, + "generated_at":"now", + "models":{"qwen3.6:27b":{"repo":"x","file":"qwen3.6-27b.mq4","size_gb":1,"min_vram_gb":1,"desc":"x"}}, + "aliases":{} + }"#; + let registry = RegistryV1::parse(raw, "test").unwrap(); + assert_eq!( + registry.resolve_tag("/elsewhere/qwen3.6-27b.mq4"), + "/elsewhere/qwen3.6-27b.mq4" + ); + assert!(registry.model("/elsewhere/qwen3.6-27b.mq4").is_none()); + // The bare file name still resolves: it carries no directory to confuse + // with the installed artifact. + assert_eq!(registry.resolve_tag("qwen3.6-27b.mq4"), "qwen3.6:27b"); + assert!(registry.model("qwen3.6-27b.mq4").is_some()); + } + + #[test] + fn entry_for_file_matches_exact_bare_names_only() { + let raw = r#"{ + "schema_version":1, + "generated_at":"now", + "models":{"qwen3.6:27b":{"repo":"x","file":"qwen3.6-27b.mq4","size_gb":1,"min_vram_gb":1,"desc":"x"}}, + "aliases":{"qwen36":"qwen3.6:27b"} + }"#; + let registry = RegistryV1::parse(raw, "test").unwrap(); + let (tag, _) = registry.entry_for_file("qwen3.6-27b.mq4").unwrap(); + assert_eq!(tag, "qwen3.6:27b"); + assert!(registry.entry_for_file("/elsewhere/qwen3.6-27b.mq4").is_none()); + assert!(registry.entry_for_file("qwen3.6:27b").is_none()); + assert!(registry.entry_for_file("qwen36").is_none()); + assert!(registry.entry_for_file("other.mq4").is_none()); + } + + #[test] fn recommended_settings_lower_the_full_sampling_contract_to_config() { let settings = RecommendedSettings { diff --git a/crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs b/crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs index 4a10c0461..3fbf3318e 100644 --- a/crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs +++ b/crates/hipfire-runtime/examples/mq4v2_gemm_parity.rs @@ -1,4 +1,5 @@ -//! v1-vs-v2 cross-check for the **WMMA GEMM** path (qt=13 vs qt=44). +//! v1-vs-v2 cross-check for the **WMMA GEMM** path (qt=13 vs qt=44), +//! plus a discriminating disjoint-halves arm for the v2 residual path. //! //! `mq4v2_parity` verifies the decode GEMV against a host oracle. It does NOT //! cover the WMMA prefill GEMMs, which are what `--scoring-mode prefill` actually @@ -6,7 +7,27 @@ //! executing (8 v2 modules compiled), WT2 KLD came back 16.705139 against a //! 0.043776 baseline. //! -//! ## Why cross-check instead of a host reference +//! ## Arm 1: Gaussian v1-vs-v2 agreement — CANNOT catch a half-select bug +//! +//! Realistic post-FWHT weights (Gaussian, sigma ~0.011) give the two halves of +//! every group near-identical `(scale, zero)` headers, so a wrong half-select +//! predicate lands inside 4-bit quantization noise. This arm detects gross v2 +//! decode errors (agreement far above the quantization floor) but a wrong +//! predicate passes it silently. See arm 2 for the discriminating fixture. +//! +//! ## Arm 2: disjoint halves + negative control — CATCHES a half-select bug +//! +//! Same construction as `mq4v2_residual_parity.rs`: half 0 in `[-1, 1]`, half 1 +//! in `[96, 160]`, packed through the same fp16 round-trip. A kernel that +//! decodes half 1 with half 0's header reconstructs `~0` instead of `~128`, +//! so the v2 output is asserted against an exact-dequant f32 reference within +//! a tight tolerance (rel-RMS below 5%). The negative control — a reference +//! computed with the halves' headers swapped — must DISAGREE by an order of +//! magnitude more; if it ever agrees, the fixture has stopped separating the +//! halves and the arm is vacuous (same control as +//! `rdna-compute/examples/mq4v2_moe_parity.rs`). +//! +//! ## Why cross-check instead of a host reference (arm 1) //! //! Replicating a WMMA kernel on the host means reproducing fp16 activation //! conversion, 16x16 tiling, and accumulation order — a reference that is itself @@ -16,10 +37,11 @@ //! Both paths then share every stage except the 8 header bytes and their decode. //! v1 quantizes with one affine grid per 256 weights; v2 with one per 128. v2 is //! therefore slightly MORE accurate, so agreement should sit at the scale of -//! 4-bit quantization noise. A systematic blow-up isolates the v2 header decode — -//! in practice the half-select predicate, which the spec calls out as "the single -//! highest-risk detail in the port" because a wrong one "compiles, runs, and -//! silently applies the wrong scale to half of every tensor." +//! 4-bit quantization noise. A systematic blow-up isolates a gross v2 header +//! decode error — but NOT the half-select predicate, which the spec calls out as +//! "the single highest-risk detail in the port" because a wrong one "compiles, +//! runs, and silently applies the wrong scale to half of every tensor." Only +//! arm 2's disjoint fixture can see that failure. //! //! ## Why sweep batch size //! @@ -27,7 +49,7 @@ //! body by batch size and flags. Scoring compiled `_bt8` and `_bt12`, so the BT //! bodies are live — and BT is b-transposed, which changes the nibble addressing //! the half-select must be derived from. Sweeping batch size tells us WHICH body -//! is wrong rather than just that something is. +//! is wrong rather than just that something is. Both arms sweep. //! //! Run: `cargo run --release -p hipfire-runtime --example mq4v2_gemm_parity` @@ -51,6 +73,7 @@ fn prng(i: usize, salt: u32) -> f32 { /// Realistic post-FWHT weights: roughly Gaussian, sigma ~0.011 as measured on the /// Qwen3.8-27B parent. Deliberately NOT the disjoint-halves fixture -- here both /// containers must be individually reasonable so their outputs are comparable. +/// (The discriminating fixture is `build_disjoint_halves` below, used by arm 2.) fn build_weights(m: usize, k: usize) -> Vec { let mut w = vec![0.0f32; m * k]; for (i, v) in w.iter_mut().enumerate() { @@ -62,6 +85,43 @@ fn build_weights(m: usize, k: usize) -> Vec { w } +/// Discriminating fixture (same construction as `mq4v2_residual_parity.rs`): +/// half 0 in `[-1, 1]`, half 1 in `[96, 160]`. The two halves occupy disjoint +/// ranges, so a kernel that decodes half 1 with half 0's header reconstructs +/// `~0` instead of `~128` and fails by >100% relative error instead of hiding +/// inside quantization noise. +fn build_disjoint_halves(m: usize, k: usize) -> Vec { + let mut w = vec![0.0f32; m * k]; + for r in 0..m { + for c in 0..k { + let gi = c % GROUP; + let idx = r * k + c; + if gi < HALF { + // [-1, 1] + w[idx] = prng(idx, 0xA5A5_0001) * 2.0 - 1.0; + } else { + // [96, 160] — disjoint from half0 by two orders of magnitude + w[idx] = 96.0 + prng(idx, 0x5A5A_0002) * 64.0; + } + } + } + w +} + +/// Swap the two 4-byte half-headers of every group in a packed v2 blob. The +/// result decodes each half with the OTHER half's grid — the negative control: +/// a reference built from this blob must DISAGREE with the correct reference. +fn swap_v2_half_headers(blob: &[u8]) -> Vec { + let mut out = blob.to_vec(); + for chunk in out.chunks_exact_mut(GROUP_BYTES) { + let mut tmp = [0u8; 8]; + tmp.copy_from_slice(&chunk[0..8]); + chunk[0..4].copy_from_slice(&tmp[4..8]); + chunk[4..8].copy_from_slice(&tmp[0..4]); + } + out +} + /// qt=13 / HFQ4 container: `[0..4) f32 scale, [4..8) f32 zero` over all 256. fn pack_v1(w: &[f32], m: usize, k: usize) -> Vec { let gpr = k / GROUP; @@ -250,6 +310,57 @@ fn main() { } } + // ── Arm 2: disjoint halves + negative control (residual path) ────────── + // + // Arm 1's Gaussian weights cannot discriminate a wrong half-select; this + // arm can. Same batch-size sweep, same v2 residual WMMA kernel, but the + // weights put half 0 in [-1, 1] and half 1 in [96, 160], so decoding half + // 1 with half 0's header is a ~100x scale error. The v2 output must match + // the exact-dequant f32 reference within a tight tolerance, AND the + // swapped-headers reference must DISAGREE — otherwise the fixture is + // vacuous and the arm proves nothing. + { + let wd = build_disjoint_halves(m, k); + let bd = pack_v2(&wd, m, k); + let bd_swapped = swap_v2_half_headers(&bd); + for &batch in &[1usize, 8, 12, 16, 32] { + let x: Vec = (0..batch * k) + .map(|i| prng(i, 0xC0FF_EE00) * 2.0 - 1.0) + .collect(); + let want = ref_gemm(&bd, &x, m, k, batch, true); + let want_bug = ref_gemm(&bd_swapped, &x, m, k, batch, true); + let bug_rel = { + let bug_f32: Vec = want_bug.iter().map(|&v| v as f32).collect(); + rel_rms(&bug_f32, &want) + }; + // Host-side negative control: the swapped grid must be badly wrong + // before any GPU result is scored against it. + assert!( + bug_rel > 0.5, + "disjoint fixture not discriminating at batch {batch}: bug_rel {bug_rel:.3e} — halves overlap" + ); + let d_a = gpu.upload_raw(&bd, &[bd.len()]).unwrap(); + let d_x = gpu.upload_f32(&x, &[batch * k]).unwrap(); + let d_y = gpu.zeros(&[batch * m], rdna_compute::DType::F32).unwrap(); + gpu.gemm_hfq4g256_residual_wmma_gfx12_mq4v2(&d_a, &d_x, &d_y, m, k, batch) + .expect("v2 disjoint residual wmma launch"); + gpu.hip.device_synchronize().unwrap(); + let got = gpu.download_f32(&d_y).unwrap(); + let e = rel_rms(&got, &want); + let verdict = if e < 0.05 && e < bug_rel * 0.1 { + "ok" + } else { + "FAIL" + }; + eprintln!( + "disjoint batch {batch:>3}: v2 rel-rms {e:.4e} bug {bug_rel:.3e} {verdict}" + ); + if verdict == "FAIL" { + failures.push((3000 + batch, e, bug_rel)); + } + } + } + // ── The fused multi-output GEMMs ──────────────────────────────────────── // // These are the rest of the live v2 set. `gemm_qkvza` carries NINETEEN header @@ -356,19 +467,63 @@ fn main() { if failures.is_empty() { eprintln!( - "\nmq4v2_gemm_parity: PASS — every live v2 WMMA GEMM matches its own exact dequant" + "\nmq4v2_gemm_parity: PASS — every live v2 WMMA GEMM matches its own exact dequant, and the disjoint-halves residual arm is half-select correct" ); } else { eprintln!( "\nmq4v2_gemm_parity: FAIL — codes {:?}", failures.iter().map(|f| f.0).collect::>() ); - eprintln!("(1000 = gate_up, 2000 = qkvza, otherwise the residual batch size)"); + eprintln!("(1000 = gate_up, 2000 = qkvza, 3000+batch = disjoint-halves residual, otherwise the Gaussian residual batch size)"); eprintln!("The v1 row is the WMMA fp16 error floor; a v2 row far above it means that"); - eprintln!("kernel mis-decodes its own header. Each body has its OWN nibble addressing,"); - eprintln!( - "so its half-select predicate must be derived from that addressing, never copied." - ); + eprintln!("kernel mis-decodes its own header. Only the disjoint arm (3000+batch) can"); + eprintln!("see a half-select bug — the Gaussian arm hides one inside quantization noise."); + eprintln!("Each body has its OWN nibble addressing, so its half-select predicate must"); + eprintln!("be derived from that addressing, never copied."); std::process::exit(1); } } + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn disjoint_fixture_discriminates_half_select() { + // Host-side proof that arm 2's negative control is load-bearing: the + // halves occupy disjoint ranges, header-swapping is an involution, + // and a swapped-headers reference DISAGREES with the correct one. + // Runs with no GPU. + let (m, k, batch) = (16usize, 256usize, 4usize); + let w = build_disjoint_halves(m, k); + for r in 0..m { + for c in 0..k { + let v = w[r * k + c]; + if (c % GROUP) < HALF { + assert!((-1.0..=1.0).contains(&v), "half0 out of range: {v}"); + } else { + assert!((96.0..=160.0).contains(&v), "half1 out of range: {v}"); + } + } + } + let blob = pack_v2(&w, m, k); + let swapped = swap_v2_half_headers(&blob); + assert_ne!(swapped, blob, "swapping identical headers would be vacuous"); + assert_eq!( + swap_v2_half_headers(&swapped), + blob, + "header swap must be an involution" + ); + let x: Vec = (0..batch * k) + .map(|i| prng(i, 0xC0FF_EE00) * 2.0 - 1.0) + .collect(); + let want = ref_gemm(&blob, &x, m, k, batch, true); + let want_bug = ref_gemm(&swapped, &x, m, k, batch, true); + let bug_f32: Vec = want_bug.iter().map(|&v| v as f32).collect(); + let bug_rel = rel_rms(&bug_f32, &want); + assert!( + bug_rel > 0.5, + "swapped-headers reference must DISAGREE: bug_rel {bug_rel:.3e}" + ); + } +} diff --git a/crates/hipfire-runtime/map.md b/crates/hipfire-runtime/map.md index 1995fd991..5c254e4dc 100644 --- a/crates/hipfire-runtime/map.md +++ b/crates/hipfire-runtime/map.md @@ -38,8 +38,8 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | [`src/cpu_router.rs`](src/cpu_router.rs) | 200 | 4 | 3 | | [`src/ddtree.rs`](src/ddtree.rs) | 2,046 | 17 | 24 | | [`src/device_mesh.rs`](src/device_mesh.rs) | 582 | 21 | 8 | -| [`src/dflash.rs`](src/dflash.rs) | 3,460 | 44 | 4 | -| [`src/dflash_generic.rs`](src/dflash_generic.rs) | 1,365 | 3 | 13 | +| [`src/dflash.rs`](src/dflash.rs) | 3,657 | 44 | 4 | +| [`src/dflash_generic.rs`](src/dflash_generic.rs) | 1,382 | 3 | 13 | | [`src/dspark_block_controller.rs`](src/dspark_block_controller.rs) | 442 | 0 | 10 | | [`src/dspark_core.rs`](src/dspark_core.rs) | 1,773 | 11 | 0 | | [`src/emit_text.rs`](src/emit_text.rs) | 1,702 | 20 | 39 | @@ -53,13 +53,13 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | [`src/kv_backend.rs`](src/kv_backend.rs) | 129 | 1 | 7 | | [`src/kv_mode.rs`](src/kv_mode.rs) | 298 | 10 | 7 | | [`src/lib.rs`](src/lib.rs) | 80 | 55 | 0 | -| [`src/llama.rs`](src/llama.rs) | 8,738 | 83 | 42 | +| [`src/llama.rs`](src/llama.rs) | 8,792 | 85 | 42 | | [`src/llama_spec.rs`](src/llama_spec.rs) | 617 | 6 | 1 | -| [`src/loader_api.rs`](src/loader_api.rs) | 256 | 10 | 4 | +| [`src/loader_api.rs`](src/loader_api.rs) | 261 | 10 | 4 | | [`src/loop_guard.rs`](src/loop_guard.rs) | 194 | 8 | 4 | | [`src/model_load.rs`](src/model_load.rs) | 117 | 8 | 1 | | [`src/model_source.rs`](src/model_source.rs) | 92 | 4 | 0 | -| [`src/multi_gpu.rs`](src/multi_gpu.rs) | 2,141 | 36 | 9 | +| [`src/multi_gpu.rs`](src/multi_gpu.rs) | 2,207 | 37 | 10 | | [`src/ngram_mod.rs`](src/ngram_mod.rs) | 484 | 11 | 13 | | [`src/paro.rs`](src/paro.rs) | 424 | 9 | 3 | | [`src/prefix.rs`](src/prefix.rs) | 109 | 3 | 6 | @@ -114,13 +114,13 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside - [`src/kv_backend.rs`](src/kv_backend.rs): `saddle_core` - [`src/kv_mode.rs`](src/kv_mode.rs): `saddle_core`, `KvModePolicy`, `ResolveResult`, `QWEN35_HFQ_POLICY`, `QWEN35_PARO_POLICY`, `DIR_SAFETENSORS_POLICY`, `LLAMA_HFQ_POLICY`, `HFQ_Q8_ONLY_POLICY`, `QWEN35_PP_POLICY`, `resolve` - [`src/lib.rs`](src/lib.rs): `admission`, `arch`, `arch_mapping`, `arch_model`, `arch_spec`, `augmentor`, `bf16_loader`, `cache_plan`, `cask`, `config`, `cpu_router`, `ddtree`, +43 more -- [`src/llama.rs`](src/llama.rs): `ModelArch`, `LlamaConfig`, `from_gguf`, `dequantize_q4_0`, `dequantize_q8_0`, `f16_to_f32`, `f32_to_f16`, `dequantize_q4_k`, `convert_q4k_to_q4f16_g64`, `convert_q4k_to_q4f16_g32`, `dequantize_q6_k`, `ParoRotation`, +71 more +- [`src/llama.rs`](src/llama.rs): `ModelArch`, `LlamaConfig`, `from_gguf`, `dequantize_q4_0`, `dequantize_q8_0`, `f16_to_f32`, `f32_to_f16`, `dequantize_q4_k`, `convert_q4k_to_q4f16_g64`, `convert_q4k_to_q4f16_g32`, `dequantize_q6_k`, `ParoRotation`, +73 more - [`src/llama_spec.rs`](src/llama_spec.rs): `verify_block_argmax`, `verify_block_logits`, `verify_block_argmax_capture_gpu`, `verify_block_sampled_capture_gpu`, `verify_tree_logits`, `lm_head_logits_n_rows` - [`src/loader_api.rs`](src/loader_api.rs): `ModelSource`, `from_path`, `arch_id`, `is_dir`, `describe`, `LoadCtx`, `SpecLoadCfg`, `CaskConfig`, `physical_cap`, `physical_cap_with_override` - [`src/loop_guard.rs`](src/loop_guard.rs): `StopReason`, `LoopGuard`, `from_config`, `new`, `off`, `enabled`, `check`, `window_len` - [`src/model_load.rs`](src/model_load.rs): `Layout`, `single`, `from_gpus`, `device_for_layer`, `output_device`, `LoadedWeights`, `WeightSource`, `load_weights` - [`src/model_source.rs`](src/model_source.rs): `TensorInfo`, `QuantConfig`, `ModelSource`, `open_model` -- [`src/multi_gpu.rs`](src/multi_gpu.rs): `BoundaryEvent`, `PeerReduceScratchLease`, `peer_reduce_scratch_bytes_per_rank`, `peer_reduce_scratch_total_bytes`, `Gpus`, `init_uniform`, `init_layers`, `init_vram_weighted`, `single`, `init_tp`, `can_access_peer_all`, `enable_peer_all`, +24 more +- [`src/multi_gpu.rs`](src/multi_gpu.rs): `BoundaryEvent`, `PeerReduceScratchLease`, `peer_reduce_scratch_bytes_per_rank`, `peer_reduce_scratch_total_bytes`, `Gpus`, `init_uniform`, `init_layers`, `init_vram_weighted`, `single`, `init_tp`, `init_ep`, `can_access_peer_all`, +25 more - [`src/ngram_mod.rs`](src/ngram_mod.rs): `HASH_MUL`, `EMPTY`, `NgramModConfig`, `NgramModPool`, `new`, `config`, `occupied`, `clear`, `insert_range`, `draft`, `record_draft_result` - [`src/paro.rs`](src/paro.rs): `repack_awq_to_hfq4g128`, `paro_text_prefix`, `load_paro_weight`, `paro_load_wt`, `paro_load_norm`, `paro_load_f32`, `alias_paro_rotation`, `load_fp16_weight_from_source`, `paro_repack_moe_projection` - [`src/prefix.rs`](src/prefix.rs): `lcp`, `TurnPlan`, `plan_turn` @@ -156,6 +156,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 58 modules · 53,079 lines · 893 public items · 617 tests · 132 examples +- 58 modules · 53,418 lines · 896 public items · 618 tests · 132 examples diff --git a/crates/hipfire-runtime/src/dflash.rs b/crates/hipfire-runtime/src/dflash.rs index 1f7ce9ad3..3d6db4ba0 100644 --- a/crates/hipfire-runtime/src/dflash.rs +++ b/crates/hipfire-runtime/src/dflash.rs @@ -617,6 +617,39 @@ fn hfq_weight( Ok(wt) } +impl DflashLayerWeights { + /// Consume one layer's weights, releasing every GPU tensor including the + /// AWQ/paro sidecars (`free_all`, not `.buf` — an AWQ-trunk drafter + /// carries one scale tensor per weight per layer). Shared by the + /// success-path `DflashWeights::free_gpu` and the `load` error path so a + /// failed load frees completed layers without duplicating this list. + fn free_gpu(self, gpu: &mut Gpu) { + let _ = gpu.free_tensor(self.attn_norm); + self.wq.free_all(gpu); + self.wk.free_all(gpu); + self.wv.free_all(gpu); + self.wo.free_all(gpu); + let _ = gpu.free_tensor(self.q_norm); + let _ = gpu.free_tensor(self.k_norm); + let _ = gpu.free_tensor(self.ffn_norm); + self.w_gate.free_all(gpu); + self.w_up.free_all(gpu); + self.w_down.free_all(gpu); + if let Some(t) = self.attn_conv_base { + let _ = gpu.free_tensor(t); + } + if let Some(w) = self.attn_conv_proj { + w.free_all(gpu); + } + if let Some(t) = self.mlp_conv_base { + let _ = gpu.free_tensor(t); + } + if let Some(w) = self.mlp_conv_proj { + w.free_all(gpu); + } + } +} + impl DflashWeights { /// True when the selector (candidate proposal) path is available. pub fn has_candidate_selector(&self) -> bool { @@ -625,22 +658,105 @@ impl DflashWeights { && self.successor_codebook.is_some() } pub fn load(gpu: &mut Gpu, hfq: &HfqFile, cfg: &DflashConfig) -> HipResult { - let fc = hfq_weight( + // Transactional construction. Every GPU tensor allocated below is + // recorded in `live_t` (plain F32) / `live_w` (weight + sidecars) and + // taken exactly once into its final owner; completed layers accumulate + // in `layers`. Any failure frees the completed layers plus every + // recorded-but-unplaced tensor before returning Err — a bare `?` + // would leak them (`GpuTensor`/`DeviceBuffer` have no `Drop`). The + // leaf loaders (`hfq_weight`, `hfq_tensor_f32`) are single-alloc and + // need no cover. Mirrors the `or_free!` style in + // `hipfire-arch-qwen35`'s `load_dflash_state`. + let mut live_t: Vec> = Vec::new(); + let mut live_w: Vec> = Vec::new(); + let mut layers: Vec = Vec::with_capacity(cfg.n_layers); + macro_rules! gt { + ($e:expr) => {{ + match $e { + Ok(t) => { + live_t.push(Some(t)); + live_t.len() - 1 + } + Err(e) => { + for l in layers.drain(..) { + l.free_gpu(gpu); + } + for slot in live_w.iter_mut() { + if let Some(w) = slot.take() { + w.free_all(gpu); + } + } + for slot in live_t.iter_mut() { + if let Some(t) = slot.take() { + let _ = gpu.free_tensor(t); + } + } + return Err(e); + } + } + }}; + } + macro_rules! wt { + ($e:expr) => {{ + match $e { + Ok(w) => { + live_w.push(Some(w)); + live_w.len() - 1 + } + Err(e) => { + for l in layers.drain(..) { + l.free_gpu(gpu); + } + for slot in live_w.iter_mut() { + if let Some(w) = slot.take() { + w.free_all(gpu); + } + } + for slot in live_t.iter_mut() { + if let Some(t) = slot.take() { + let _ = gpu.free_tensor(t); + } + } + return Err(e); + } + } + }}; + } + macro_rules! take_t { + ($i:expr) => { + live_t[$i] + .take() + .expect("dflash load: F32 slot taken twice") + }; + } + macro_rules! take_w { + ($i:expr) => { + live_w[$i] + .take() + .expect("dflash load: weight slot taken twice") + }; + } + let i_fc = wt!(hfq_weight( hfq, gpu, "fc.weight", cfg.hidden, cfg.num_extract() * cfg.hidden, - )?; - let hidden_norm = hfq_tensor_f32(hfq, gpu, "hidden_norm.weight", vec![cfg.hidden])?; - let norm = hfq_tensor_f32(hfq, gpu, "norm.weight", vec![cfg.hidden])?; + )); + let i_hidden_norm = gt!(hfq_tensor_f32( + hfq, + gpu, + "hidden_norm.weight", + vec![cfg.hidden] + )); + let i_norm = gt!(hfq_tensor_f32(hfq, gpu, "norm.weight", vec![cfg.hidden])); let conv_k = cfg.conv_kernel_size.unwrap_or(2); let conv_g = cfg.conv_group_size.unwrap_or(16); let conv_groups = cfg.hidden / conv_g; let proj_m = 2 * conv_k * conv_groups; - let mut layers = Vec::with_capacity(cfg.n_layers); + // (`layers` is declared above so the `gt!`/`wt!` error arms can free it.) for i in 0..cfg.n_layers { let p = format!("layers.{i}"); // Attempt DFlash2 conv weights; absent on legacy drafts. @@ -656,12 +772,12 @@ impl DflashWeights { }; if hfq.tensor_data(&key).is_some() { // shape 2*K*H - Some(hfq_tensor_f32( + Some(gt!(hfq_tensor_f32( hfq, gpu, &key, vec![2 * conv_k * cfg.hidden], - )?) + ))) } else { None } @@ -677,7 +793,7 @@ impl DflashWeights { alt }; if hfq.tensor_data(&key).is_some() { - Some(hfq_weight(hfq, gpu, &key, proj_m, cfg.hidden)?) + Some(wt!(hfq_weight(hfq, gpu, &key, proj_m, cfg.hidden))) } else { None } @@ -693,12 +809,12 @@ impl DflashWeights { alt }; if hfq.tensor_data(&key).is_some() { - Some(hfq_tensor_f32( + Some(gt!(hfq_tensor_f32( hfq, gpu, &key, vec![2 * conv_k * cfg.hidden], - )?) + ))) } else { None } @@ -714,110 +830,125 @@ impl DflashWeights { alt }; if hfq.tensor_data(&key).is_some() { - Some(hfq_weight(hfq, gpu, &key, proj_m, cfg.hidden)?) + Some(wt!(hfq_weight(hfq, gpu, &key, proj_m, cfg.hidden))) } else { None } } else { None }; - let layer = DflashLayerWeights { - attn_norm: hfq_tensor_f32( - hfq, - gpu, - &format!("{p}.input_layernorm.weight"), - vec![cfg.hidden], - )?, - wq: hfq_weight( - hfq, - gpu, - &format!("{p}.self_attn.q_proj.weight"), - cfg.q_dim(), - cfg.hidden, - )?, - wk: hfq_weight( - hfq, - gpu, - &format!("{p}.self_attn.k_proj.weight"), - cfg.kv_dim(), - cfg.hidden, - )?, - wv: hfq_weight( - hfq, - gpu, - &format!("{p}.self_attn.v_proj.weight"), - cfg.kv_dim(), - cfg.hidden, - )?, - wo: hfq_weight( - hfq, - gpu, - &format!("{p}.self_attn.o_proj.weight"), - cfg.hidden, - cfg.q_dim(), - )?, - q_norm: hfq_tensor_f32( - hfq, - gpu, - &format!("{p}.self_attn.q_norm.weight"), - vec![cfg.head_dim], - )?, - k_norm: hfq_tensor_f32( - hfq, - gpu, - &format!("{p}.self_attn.k_norm.weight"), - vec![cfg.head_dim], - )?, - ffn_norm: hfq_tensor_f32( - hfq, - gpu, - &format!("{p}.post_attention_layernorm.weight"), - vec![cfg.hidden], - )?, - w_gate: hfq_weight( - hfq, - gpu, - &format!("{p}.mlp.gate_proj.weight"), - cfg.intermediate, - cfg.hidden, - )?, - w_up: hfq_weight( - hfq, - gpu, - &format!("{p}.mlp.up_proj.weight"), - cfg.intermediate, - cfg.hidden, - )?, - w_down: hfq_weight( - hfq, - gpu, - &format!("{p}.mlp.down_proj.weight"), - cfg.hidden, - cfg.intermediate, - )?, - attn_conv_base, - attn_conv_proj, - mlp_conv_base, - mlp_conv_proj, - }; - layers.push(layer); + // Stage every field as a slot index first: a failure below frees + // the staged slots (plus completed layers) via `gt!`/`wt!`, and + // the `take_*!` push itself is infallible. + let i_attn_norm = gt!(hfq_tensor_f32( + hfq, + gpu, + &format!("{p}.input_layernorm.weight"), + vec![cfg.hidden], + )); + let i_wq = wt!(hfq_weight( + hfq, + gpu, + &format!("{p}.self_attn.q_proj.weight"), + cfg.q_dim(), + cfg.hidden, + )); + let i_wk = wt!(hfq_weight( + hfq, + gpu, + &format!("{p}.self_attn.k_proj.weight"), + cfg.kv_dim(), + cfg.hidden, + )); + let i_wv = wt!(hfq_weight( + hfq, + gpu, + &format!("{p}.self_attn.v_proj.weight"), + cfg.kv_dim(), + cfg.hidden, + )); + let i_wo = wt!(hfq_weight( + hfq, + gpu, + &format!("{p}.self_attn.o_proj.weight"), + cfg.hidden, + cfg.q_dim(), + )); + let i_q_norm = gt!(hfq_tensor_f32( + hfq, + gpu, + &format!("{p}.self_attn.q_norm.weight"), + vec![cfg.head_dim], + )); + let i_k_norm = gt!(hfq_tensor_f32( + hfq, + gpu, + &format!("{p}.self_attn.k_norm.weight"), + vec![cfg.head_dim], + )); + let i_ffn_norm = gt!(hfq_tensor_f32( + hfq, + gpu, + &format!("{p}.post_attention_layernorm.weight"), + vec![cfg.hidden], + )); + let i_w_gate = wt!(hfq_weight( + hfq, + gpu, + &format!("{p}.mlp.gate_proj.weight"), + cfg.intermediate, + cfg.hidden, + )); + let i_w_up = wt!(hfq_weight( + hfq, + gpu, + &format!("{p}.mlp.up_proj.weight"), + cfg.intermediate, + cfg.hidden, + )); + let i_w_down = wt!(hfq_weight( + hfq, + gpu, + &format!("{p}.mlp.down_proj.weight"), + cfg.hidden, + cfg.intermediate, + )); + layers.push(DflashLayerWeights { + attn_norm: take_t!(i_attn_norm), + wq: take_w!(i_wq), + wk: take_w!(i_wk), + wv: take_w!(i_wv), + wo: take_w!(i_wo), + q_norm: take_t!(i_q_norm), + k_norm: take_t!(i_k_norm), + ffn_norm: take_t!(i_ffn_norm), + w_gate: take_w!(i_w_gate), + w_up: take_w!(i_w_up), + w_down: take_w!(i_w_down), + attn_conv_base: attn_conv_base.map(|j| take_t!(j)), + attn_conv_proj: attn_conv_proj.map(|j| take_w!(j)), + mlp_conv_base: mlp_conv_base.map(|j| take_t!(j)), + mlp_conv_proj: mlp_conv_proj.map(|j| take_w!(j)), + }); } // Selector: hidden_projection [rank, hidden] + two codebooks [vocab, rank] host-side // Exact HFQ names are `candidate_selector.hidden_projection.weight`, // `candidate_selector.predecessor_codebook`, `candidate_selector.successor_codebook`. // Optional fallback `.weight` suffix is tolerated but not required. - let selector_hidden_proj = if cfg.selector_rank.is_some() { + // Slot index (`take_w!`n below); the codebooks between here and the + // take are host-side only and cannot fail with GPU memory held. + let i_selector_hidden_proj = if cfg.selector_rank.is_some() { let rank = cfg.selector_rank.unwrap(); let candidates = [ "candidate_selector.hidden_projection.weight", "selector.hidden_projection.weight", "selector.hidden_proj.weight", ]; - let mut found = None; + let mut found: Option = None; for n in candidates { if hfq.tensor_data(n).is_some() { - found = Some(hfq_weight(hfq, gpu, n, rank, cfg.hidden)?); + found = Some(wt!(hfq_weight(hfq, gpu, n, rank, cfg.hidden))); break; } } @@ -896,6 +1027,12 @@ impl DflashWeights { (None, None, None, None) }; + let fc = take_w!(i_fc); + let hidden_norm = take_t!(i_hidden_norm); + let norm = take_t!(i_norm); + let selector_hidden_proj = i_selector_hidden_proj.map(|j| take_w!(j)); + debug_assert!(live_t.iter().all(|s| s.is_none())); + debug_assert!(live_w.iter().all(|s| s.is_none())); let has_mq = std::iter::once(&fc) .chain(layers.iter().flat_map(|l| { let mut v: Vec<&WeightTensor> = @@ -924,8 +1061,20 @@ impl DflashWeights { }); if has_mq { // MQ dispatch needs the engine's FWHT sign tables uploaded - // (matches `gemv_mq4g256_with_rotate`'s setup). - gpu.ensure_mq_signs()?; + // (matches `gemv_mq4g256_with_rotate`'s setup). A failure here + // must still release the weights above before returning Err. + if let Err(e) = gpu.ensure_mq_signs() { + fc.free_all(gpu); + let _ = gpu.free_tensor(hidden_norm); + let _ = gpu.free_tensor(norm); + for l in layers { + l.free_gpu(gpu); + } + if let Some(w) = selector_hidden_proj { + w.free_all(gpu); + } + return Err(e); + } } Ok(DflashWeights { @@ -951,29 +1100,7 @@ impl DflashWeights { let _ = gpu.free_tensor(self.hidden_norm); let _ = gpu.free_tensor(self.norm); for l in self.layers { - let _ = gpu.free_tensor(l.attn_norm); - l.wq.free_all(gpu); - l.wk.free_all(gpu); - l.wv.free_all(gpu); - l.wo.free_all(gpu); - let _ = gpu.free_tensor(l.q_norm); - let _ = gpu.free_tensor(l.k_norm); - let _ = gpu.free_tensor(l.ffn_norm); - l.w_gate.free_all(gpu); - l.w_up.free_all(gpu); - l.w_down.free_all(gpu); - if let Some(t) = l.attn_conv_base { - let _ = gpu.free_tensor(t); - } - if let Some(w) = l.attn_conv_proj { - w.free_all(gpu); - } - if let Some(t) = l.mlp_conv_base { - let _ = gpu.free_tensor(t); - } - if let Some(w) = l.mlp_conv_proj { - w.free_all(gpu); - } + l.free_gpu(gpu); } if let Some(w) = self.selector_hidden_proj { w.free_all(gpu); @@ -1346,13 +1473,36 @@ impl DflashScratch { if let Some(v) = s.v_ctx_cached.pop() { let _ = gpu.free_tensor(v); } - s.k_full_cached = Some(gpu.alloc_tensor(&[w_full * kvd], DType::F32)?); - s.v_full_cached = Some(gpu.alloc_tensor(&[w_full * kvd], DType::F32)?); - s.k_cat_full = Some(gpu.alloc_tensor(&[(w_full + b) * kvd], DType::F32)?); - s.v_cat_full = Some(gpu.alloc_tensor(&[(w_full + b) * kvd], DType::F32)?); + // The base scratch `s` is fully owned here: any failure below frees + // it before returning Err — a bare `?` would leak it (no `Drop` on + // the GPU-owning types), including when the failure follows the pop + // above. Same class as the `or_free!` sites in `load_dflash_state`. + // + // Each allocation is parked in `s` the moment it succeeds, so the + // error arm's `s.free_gpu` also covers every earlier allocation of + // this ladder: with the tensors held as locals until the end, a + // failure on the 2nd..5th alloc freed `s` but leaked the locals + // (hw-gate Fable seat on #691, run 33900101473). + macro_rules! alloc_or_free { + ($e:expr) => { + match $e { + Ok(t) => t, + Err(e) => { + s.free_gpu(gpu); + return Err(e); + } + } + }; + } + s.k_full_cached = Some(alloc_or_free!(gpu.alloc_tensor(&[w_full * kvd], DType::F32))); + s.v_full_cached = Some(alloc_or_free!(gpu.alloc_tensor(&[w_full * kvd], DType::F32))); + s.k_cat_full = Some(alloc_or_free!(gpu.alloc_tensor(&[(w_full + b) * kvd], DType::F32))); + s.v_cat_full = Some(alloc_or_free!(gpu.alloc_tensor(&[(w_full + b) * kvd], DType::F32))); // positions_k holds the last w_full context rows + the B noise rows // (the forward uploads only that suffix; every layer's span is one). - let new_positions_k = gpu.alloc_tensor(&[w_full + b], DType::F32)?; + // Allocate before freeing the old buffer so a failure still leaves + // `s` intact for the error arm above. + let new_positions_k = alloc_or_free!(gpu.alloc_tensor(&[w_full + b], DType::F32)); let _ = gpu.free_tensor(std::mem::replace(&mut s.positions_k, new_positions_k)); // The ctx bound is the target's physical capacity, not the window — // l may cross w_full (the last layer's span just slides). @@ -1380,7 +1530,38 @@ impl DflashScratch { let qd = cfg.q_dim(); let kvd = cfg.kv_dim(); - let mq_x_rot = if with_mq { + // Transactional construction: every `alloc_tensor` below goes through + // `at!`, which records the tensor in `live`; each index is taken + // exactly once when the struct is built. On failure the error arm + // frees everything recorded so far and returns — a bare `?` would + // leak (`GpuTensor`/`DeviceBuffer` have no `Drop`). Same style as the + // `gt!`/`wt!` slots in `DflashWeights::load` above. + let mut live: Vec> = Vec::new(); + macro_rules! at { + ($shape:expr) => {{ + match gpu.alloc_tensor($shape, DType::F32) { + Ok(t) => { + live.push(Some(t)); + live.len() - 1 + } + Err(e) => { + for slot in live.iter_mut() { + if let Some(t) = slot.take() { + let _ = gpu.free_tensor(t); + } + } + return Err(e); + } + } + }}; + } + macro_rules! take { + ($i:expr) => { + live[$i].take().expect("dflash scratch slot taken twice") + }; + } + + let i_mq_x_rot = if with_mq { // Sized for a CHUNK of the worst-case MQ rotation, not the whole // first-call prefix. The rotations called through `gemm_dispatch` // are: @@ -1401,93 +1582,109 @@ impl DflashScratch { // `ceil(batch / chunk_rows)` smaller GEMMs — adds ~1-2 launches per // 1K prefix tokens (negligible vs seconds-scale prefill). let widest = MQ_X_ROT_CHUNK_ROWS * std::cmp::max(inter, std::cmp::max(qd, ne * h)); - Some(gpu.alloc_tensor(&[widest], DType::F32)?) + Some(at!(&[widest])) } else { None }; // DFlash2 optional buffers: allocated only when the config declares them. - let (conv_temp, conv_dynamic, selector_proj, topk_ids, topk_vals) = { - let need_conv = cfg.conv_kernel_size.is_some() && cfg.conv_group_size.is_some(); - let need_selector = cfg.selector_rank.is_some() && cfg.selector_top_k.is_some(); - let ct = if need_conv { - Some(gpu.alloc_tensor(&[b * h], DType::F32)?) - } else { - None - }; - let cd = if need_conv { - let k = cfg.conv_kernel_size.unwrap(); - let g = cfg.conv_group_size.unwrap(); - let groups = h / g; - let stride = 2 * k * groups; - Some(gpu.alloc_tensor(&[b * stride], DType::F32)?) - } else { - None - }; - let sp = if need_selector { - let rank = cfg.selector_rank.unwrap(); - Some(gpu.alloc_tensor(&[b * rank], DType::F32)?) - } else { - None - }; - let (ti, tv) = if need_selector { - let kk = cfg.selector_top_k.unwrap(); - // ids as i32 stored in F32 buffer (reinterprets), vals as f32 - ( - Some(gpu.alloc_tensor(&[b * kk], DType::F32)?), - Some(gpu.alloc_tensor(&[b * kk], DType::F32)?), - ) - } else { - (None, None) - }; - (ct, cd, sp, ti, tv) + // Slot indices (`take!`n at the build below). + let need_conv = cfg.conv_kernel_size.is_some() && cfg.conv_group_size.is_some(); + let need_selector = cfg.selector_rank.is_some() && cfg.selector_top_k.is_some(); + let i_conv_temp = if need_conv { Some(at!(&[b * h])) } else { None }; + let i_conv_dynamic = if need_conv { + let k = cfg.conv_kernel_size.unwrap(); + let g = cfg.conv_group_size.unwrap(); + let groups = h / g; + let stride = 2 * k * groups; + Some(at!(&[b * stride])) + } else { + None + }; + let i_selector_proj = if need_selector { + let rank = cfg.selector_rank.unwrap(); + Some(at!(&[b * rank])) + } else { + None + }; + let (i_topk_ids, i_topk_vals) = if need_selector { + let kk = cfg.selector_top_k.unwrap(); + // ids as i32 stored in F32 buffer (reinterprets), vals as f32 + (Some(at!(&[b * kk])), Some(at!(&[b * kk]))) + } else { + (None, None) }; // Per-layer cache buffers for k_ctx/v_ctx (post-norm-for-K, pre-rope). // Size each at [max_ctx × kv_dim] f32 = l × kvd × 4 bytes. Memory // cost for 16-layer / 4096-ctx / 256-kv_dim draft ≈ 2 × 16 × 4 MB // = 128 MB. Trivial vs 24 GB VRAM. - let mut k_ctx_cached = Vec::with_capacity(cfg.n_layers); - let mut v_ctx_cached = Vec::with_capacity(cfg.n_layers); + let mut kv_idx: Vec<(usize, usize)> = Vec::with_capacity(cfg.n_layers); let mut draft_ffn_graphs = Vec::with_capacity(cfg.n_layers); let mut draft_ffn_warmed_up = Vec::with_capacity(cfg.n_layers); for _ in 0..cfg.n_layers { - k_ctx_cached.push(gpu.alloc_tensor(&[l * kvd], DType::F32)?); - v_ctx_cached.push(gpu.alloc_tensor(&[l * kvd], DType::F32)?); + kv_idx.push((at!(&[l * kvd]), at!(&[l * kvd]))); draft_ffn_graphs.push(HashMap::new()); draft_ffn_warmed_up.push(HashSet::new()); } + let i_x = at!(&[b * h]); + let i_x_norm = at!(&[b * h]); + let i_q = at!(&[b * qd]); + let i_k_noise = at!(&[b * kvd]); + let i_v_noise = at!(&[b * kvd]); + let i_gate = at!(&[b * inter]); + let i_up = at!(&[b * inter]); + let i_gate_up = at!(&[b * inter]); + let i_attn_out = at!(&[b * qd]); + let i_residual = at!(&[b * h]); + + let i_target_hidden = at!(&[l * ne * h]); + let i_target_hidden_proj = at!(&[l * h]); + + let i_k_cat = at!(&[tot * kvd]); + let i_v_cat = at!(&[tot * kvd]); + + let i_positions_q = at!(&[b]); + let i_positions_k = at!(&[tot]); + + let mut k_ctx_cached = Vec::with_capacity(cfg.n_layers); + let mut v_ctx_cached = Vec::with_capacity(cfg.n_layers); + for (ik, iv) in kv_idx { + k_ctx_cached.push(take!(ik)); + v_ctx_cached.push(take!(iv)); + } + debug_assert!(live.iter().all(|s| s.is_none())); Ok(DflashScratch { max_block_size: b, max_ctx_len: l, - x: gpu.alloc_tensor(&[b * h], DType::F32)?, - x_norm: gpu.alloc_tensor(&[b * h], DType::F32)?, - q: gpu.alloc_tensor(&[b * qd], DType::F32)?, - k_noise: gpu.alloc_tensor(&[b * kvd], DType::F32)?, - v_noise: gpu.alloc_tensor(&[b * kvd], DType::F32)?, - gate: gpu.alloc_tensor(&[b * inter], DType::F32)?, - up: gpu.alloc_tensor(&[b * inter], DType::F32)?, - gate_up: gpu.alloc_tensor(&[b * inter], DType::F32)?, - attn_out: gpu.alloc_tensor(&[b * qd], DType::F32)?, - residual: gpu.alloc_tensor(&[b * h], DType::F32)?, - - target_hidden: gpu.alloc_tensor(&[l * ne * h], DType::F32)?, - target_hidden_proj: gpu.alloc_tensor(&[l * h], DType::F32)?, - - k_cat: gpu.alloc_tensor(&[tot * kvd], DType::F32)?, - v_cat: gpu.alloc_tensor(&[tot * kvd], DType::F32)?, - - positions_q: gpu.alloc_tensor(&[b], DType::F32)?, - positions_k: gpu.alloc_tensor(&[tot], DType::F32)?, - - mq_x_rot, - conv_temp, - conv_dynamic, - selector_proj, - topk_ids, - topk_vals, + x: take!(i_x), + x_norm: take!(i_x_norm), + q: take!(i_q), + k_noise: take!(i_k_noise), + v_noise: take!(i_v_noise), + gate: take!(i_gate), + up: take!(i_up), + gate_up: take!(i_gate_up), + attn_out: take!(i_attn_out), + residual: take!(i_residual), + + target_hidden: take!(i_target_hidden), + target_hidden_proj: take!(i_target_hidden_proj), + + k_cat: take!(i_k_cat), + v_cat: take!(i_v_cat), + + positions_q: take!(i_positions_q), + positions_k: take!(i_positions_k), + + mq_x_rot: i_mq_x_rot.map(|j| take!(j)), + conv_temp: i_conv_temp.map(|j| take!(j)), + conv_dynamic: i_conv_dynamic.map(|j| take!(j)), + selector_proj: i_selector_proj.map(|j| take!(j)), + topk_ids: i_topk_ids.map(|j| take!(j)), + topk_vals: i_topk_vals.map(|j| take!(j)), thlog: TargetHiddenLog::new(), k_ctx_cached, v_ctx_cached, diff --git a/crates/hipfire-runtime/src/dflash_generic.rs b/crates/hipfire-runtime/src/dflash_generic.rs index 019542907..f4c5526c0 100644 --- a/crates/hipfire-runtime/src/dflash_generic.rs +++ b/crates/hipfire-runtime/src/dflash_generic.rs @@ -1048,18 +1048,35 @@ pub fn build_generic_dflash_speculator( let block_size = config.block_size; // L3: F16 drafts (dflash_convert) → has_mq=false → DflashScratch::new. // new_with_mq only for an MQ-quantized draft. - let scratch = if weights.has_mq { + // Transactional: `weights` and `scratch` own GPU memory with no `Drop`, + // so each fallible step below frees what is already owned before + // returning Err (same class as the `or_free!` chain in + // `hipfire-arch-qwen35`'s `load_dflash_state`). + let scratch = match if weights.has_mq { DflashScratch::new_with_mq(gpu, &config, block_size, ctx_capacity, true) - .map_err(|e| format!("{e}"))? + .map_err(|e| format!("{e}")) } else { - DflashScratch::new(gpu, &config, block_size, ctx_capacity).map_err(|e| format!("{e}"))? + DflashScratch::new(gpu, &config, block_size, ctx_capacity).map_err(|e| format!("{e}")) + } { + Ok(s) => s, + Err(e) => { + weights.free_gpu(gpu); + return Err(e); + } }; let _ = draft_hfq; // Tell the target which residual-hidden layers to capture (the drafter's // target_layer_ids), and mint the per-target verify scratch. target.set_dflash_extract_layers(config.target_layer_ids.clone()); - let verify_scratch = target.new_spec_scratch(gpu, block_size)?; + let verify_scratch = match target.new_spec_scratch(gpu, block_size) { + Ok(v) => v, + Err(e) => { + weights.free_gpu(gpu); + scratch.free_gpu(gpu); + return Err(e); + } + }; Ok(Box::new(GenericDflashSpeculator { weights, diff --git a/crates/hipfire-runtime/src/llama.rs b/crates/hipfire-runtime/src/llama.rs index e9e909107..04f4203d8 100644 --- a/crates/hipfire-runtime/src/llama.rs +++ b/crates/hipfire-runtime/src/llama.rs @@ -680,18 +680,24 @@ pub struct LayerWeights { impl LlamaWeights { /// Return all GPU buffers to the pool (drained on unload). Consumes self. + /// Each weight goes through `WeightTensor::free_all` so the PARO rotation + /// and AWQ scale sidecars are released with their buffers. pub fn free_gpu(self, gpu: &mut Gpu) { let _ = gpu.free_tensor(self.token_embd); let _ = gpu.free_tensor(self.output_norm); if !self.lm_head_aliases_embd { - let _ = gpu.free_tensor(self.output.buf); + // free_all (not .buf) so the AWQ / PARO sidecars are released too. + // The tied-lm_head alias carries no sidecars by construction + // (`tied_lm_head_alias` sets paro/awq_scale to None), so skipping + // the whole output weight when aliased still frees exactly once. + self.output.free_all(gpu); } for l in self.layers { let _ = gpu.free_tensor(l.attn_norm); - let _ = gpu.free_tensor(l.wq.buf); - let _ = gpu.free_tensor(l.wk.buf); - let _ = gpu.free_tensor(l.wv.buf); - let _ = gpu.free_tensor(l.wo.buf); + l.wq.free_all(gpu); + l.wk.free_all(gpu); + l.wv.free_all(gpu); + l.wo.free_all(gpu); if let Some(t) = l.q_norm { let _ = gpu.free_tensor(t); } @@ -699,9 +705,9 @@ impl LlamaWeights { let _ = gpu.free_tensor(t); } let _ = gpu.free_tensor(l.ffn_norm); - let _ = gpu.free_tensor(l.w_gate.buf); - let _ = gpu.free_tensor(l.w_up.buf); - let _ = gpu.free_tensor(l.w_down.buf); + l.w_gate.free_all(gpu); + l.w_up.free_all(gpu); + l.w_down.free_all(gpu); } } } @@ -1838,10 +1844,50 @@ pub fn prefill_forward( /// largest physical_cap any consumer sets up. pub const PREFILL_MAX_BATCH: usize = 256; +/// Kill-switch for the MQ-V2 (qt44 + neutral qt47-50) gfx11 WMMA prefill path: +/// gfx12 (`gfx1200`/`gfx1201`) is always admitted, gfx11 +/// (`gfx1100`/`gfx1101`/`gfx1102`/`gfx1150`/`gfx1151`) is admitted unless +/// `HIPFIRE_MQV2_GFX11_WMMA=0`, anything else is rejected. Defined here as +/// the shared home for the qwen35 caller (`qwen35::is_batchable_la`); +/// `llama::is_batchable_la` does NOT delegate to it (see below). +/// `value` is the raw env var (None = unset → default ON); only `Some("0")` +/// disables the gfx11 path. Gfx12 is unaffected by the env var. +pub fn mqv2_gfx11_wmma_enabled_from_env(value: Option<&str>, arch: &str) -> bool { + let gfx11_enabled = value != Some("0"); + if matches!(arch, "gfx1200" | "gfx1201") { + true + } else if matches!( + arch, + "gfx1100" | "gfx1101" | "gfx1102" | "gfx1150" | "gfx1151" + ) { + gfx11_enabled + } else { + false + } +} + +/// Admit rule for the MQ-V2 family (`MQ4G256V2` + neutral `MQ6/5/3/2G256V2`) +/// in batched WMMA prefill: dtype set × arch set × the +/// `HIPFIRE_MQV2_GFX11_WMMA` kill-switch in one function. Only +/// `qwen35::is_batchable_la` delegates here — qwen35's +/// `forward_prefill_chunk` has V2 dispatch arms, while the llama chunk path +/// does not (see `llama::is_batchable_la`). +/// `MQ4CG256` (qt45) stays gfx12-only in its caller and is intentionally +/// NOT part of this rule. +pub fn mqv2_wmma_batchable(dt: DType, mqv2_gfx11_wmma: Option<&str>, arch: &str) -> bool { + matches!( + dt, + DType::MQ4G256V2 + | DType::MQ6G256V2 + | DType::MQ5G256V2 + | DType::MQ3G256V2 + | DType::MQ2G256V2 + ) && mqv2_gfx11_wmma_enabled_from_env(mqv2_gfx11_wmma, arch) +} + /// Is this dtype/arch combination eligible for the batched WMMA prefill -/// kernels? Matches `qwen35::is_batchable_la` exactly so plain Qwen3 and -/// hybrid Qwen3.5 share one rule and stay in lockstep when new dtypes or -/// arches gain WMMA support. +/// kernels? NOTE: unlike `qwen35::is_batchable_la`, this does NOT admit the +/// MQ-V2 family — see the `never_v2` refusal below. pub fn is_batchable_la(dt: DType, arch: &str) -> bool { let always_ok = matches!( dt, @@ -1874,21 +1920,30 @@ pub fn is_batchable_la(dt: DType, arch: &str) -> bool { arch, "gfx1010" | "gfx1011" | "gfx1012" | "gfx1013" | "gfx1030" | "gfx1031" | "gfx1032" ); - // MQ4G256V2 / MQ4CG256 batched prefill + batched lm_head GEMM exist only - // on gfx12 (gfx1200/gfx1201). Outside gfx12, fall back to per-token decode - // rather than dispatching a gfx12 WMMA kernel. Lockstep with - // qwen35::is_batchable_la (qt44/qt45). - // Extended to neutral V2 family qt47-50. - let mq4_v2_gfx12 = matches!( + // MQ-V2 family (`MQ4G256V2` + neutral `MQ6/5/3/2G256V2`, qt44/qt47-50) + // plus `MQ4CG256` (qt45): REFUSED on every arch. `forward_prefill_chunk` + // has no V2 arms — its per-layer dtype matchers (`qkv_is_mq` ~:2570, + // `wo_is_mq` ~:3025, `ffn_is_mq` ~:3117, `w_down_is_mq` ~:3248) list only + // `MQ4G256|MQ6G256|MQ3G256|MFP4G32`, so an admitted V2 model would skip + // the FWHT rotate and run the V1 `hfq4g256` launchers + // (`gemm_qkv_hfq4g256`, `gemm_hfq4g256_residual`, `gemm_gate_up_hfq4g256`) + // on V2 blobs — silently incoherent prefill. Per-token decode is the + // only correct llama path for V2 until those arms exist. qwen35's chunk + // path DOES have the V2 arms, so `qwen35::is_batchable_la` keeps + // admitting V2 via the shared `mqv2_wmma_batchable` rule above. + let never_v2 = matches!( dt, DType::MQ4G256V2 - | DType::MQ4CG256 | DType::MQ6G256V2 | DType::MQ5G256V2 | DType::MQ3G256V2 | DType::MQ2G256V2 - ) && matches!(arch, "gfx1200" | "gfx1201"); - wmma_only || mq3_gfx10_scalar || mq4_v2_gfx12 + | DType::MQ4CG256 + ); + if never_v2 { + return false; + } + wmma_only || mq3_gfx10_scalar } /// Per-call scratch for `forward_prefill_batch`. Holds [N × ...] working @@ -8247,20 +8302,18 @@ mod tests { } #[test] - fn is_batchable_la_mq4_v2_gfx12_only() { - // MQ4G256V2 / MQ4CG256 batched prefill is gfx12-only; other arches - // fall back to per-token decode. - for arch in ["gfx1200", "gfx1201"] { - assert!( - is_batchable_la(DType::MQ4G256V2, arch), - "MQ4G256V2 should batch on {arch}" - ); - assert!( - is_batchable_la(DType::MQ4CG256, arch), - "MQ4CG256 should batch on {arch}" - ); - } - for arch in ["gfx1010", "gfx1100", "gfx942"] { + fn is_batchable_la_mq4_v2_refused_everywhere() { + // `forward_prefill_chunk` has no V2 arms (its `qkv_is_mq` ~:2570, + // `wo_is_mq` ~:3025, `ffn_is_mq` ~:3117, `w_down_is_mq` ~:3248 + // matchers list only V1 dtypes), so llama must refuse MQ4G256V2 and + // MQ4CG256 on EVERY arch — including gfx11/gfx12 — and stay on + // per-token decode. qwen35's chunk path has the arms and keeps the + // shared `mqv2_wmma_batchable` rule; see + // `qwen35_is_batchable_la_mq4_v2_gfx11_and_gfx12` for the admit side. + for arch in [ + "gfx1100", "gfx1101", "gfx1102", "gfx1150", "gfx1151", "gfx1200", "gfx1201", + "gfx1010", "gfx1030", "gfx942", + ] { assert!( !is_batchable_la(DType::MQ4G256V2, arch), "MQ4G256V2 must fall back on {arch}" @@ -8273,18 +8326,19 @@ mod tests { } #[test] - fn is_batchable_la_v2_family_gfx12_only() { - for arch in ["gfx1200", "gfx1201"] { - assert!(is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 gfx12"); - assert!(is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 gfx12"); - assert!(is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 gfx12"); - assert!(is_batchable_la(DType::MQ2G256V2, arch), "MQ2V2 gfx12"); - } - for arch in ["gfx1010", "gfx1100", "gfx942"] { - assert!(!is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 fallback"); - assert!(!is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 fallback"); - assert!(!is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 fallback"); - assert!(!is_batchable_la(DType::MQ2G256V2, arch), "MQ2V2 fallback"); + fn is_batchable_la_v2_family_refused_everywhere() { + // Neutral V2 family (qt47-50): same refusal as MQ4G256V2 — no V2 arms + // in the llama chunk path, so refuse on every arch (including + // gfx11/gfx12). Mirrors `qwen35_is_batchable_la_v2_family_gfx11_and_gfx12` + // on the admit side. + for arch in [ + "gfx1100", "gfx1101", "gfx1102", "gfx1150", "gfx1151", "gfx1200", "gfx1201", + "gfx1010", "gfx1030", "gfx942", + ] { + assert!(!is_batchable_la(DType::MQ6G256V2, arch), "MQ6V2 fallback on {arch}"); + assert!(!is_batchable_la(DType::MQ5G256V2, arch), "MQ5V2 fallback on {arch}"); + assert!(!is_batchable_la(DType::MQ3G256V2, arch), "MQ3V2 fallback on {arch}"); + assert!(!is_batchable_la(DType::MQ2G256V2, arch), "MQ2V2 fallback on {arch}"); } assert_ne!(DType::MQ6G256, DType::MQ6G256V2); assert_ne!(DType::MQ3G256, DType::MQ3G256V2); diff --git a/crates/hipfire-runtime/src/loader_api.rs b/crates/hipfire-runtime/src/loader_api.rs index cb65f2cc9..5934b3200 100644 --- a/crates/hipfire-runtime/src/loader_api.rs +++ b/crates/hipfire-runtime/src/loader_api.rs @@ -132,6 +132,11 @@ pub struct SpecLoadCfg { pub mtp: Option, /// MTP draft window K. `None` = runtime default (`HIPFIRE_MTP_K`). pub mtp_k: Option, + /// Qwen DFlash draft enable, lowered from `dflash_mode`: `Some(true)` = + /// `on` (fail the load when the draft is missing or unloadable), + /// `Some(false)` = `off` (skip), `None` = `auto` (load when present, + /// log-and-AR fallback otherwise). + pub dflash: Option, } /// CASK/TriAttention params forwarded by the CLI at load time. diff --git a/crates/hipfire-runtime/src/multi_gpu.rs b/crates/hipfire-runtime/src/multi_gpu.rs index 3bc9a5492..5d05346c9 100644 --- a/crates/hipfire-runtime/src/multi_gpu.rs +++ b/crates/hipfire-runtime/src/multi_gpu.rs @@ -316,6 +316,55 @@ impl Gpus { }) } + /// Expert-parallel constructor: bring up `ep_size` devices that each run + /// **every** layer (PP=1), with routed experts sharded across ranks per an + /// expert assignment (e.g. `ExpertAssign::Stride`). + /// + /// Layout-identical to [`Gpus::init_tp`]: same device set, streams, + /// pre-flight VRAM gate, and PP=1 layer-band map (`tp_band_starts` is + /// layout-generic despite the name). Only the recorded [`DeviceMesh`] + /// axis differs — `Ep` instead of `Tp` — so `mesh.size_of(Ep)` reports + /// the rank count after an EP load instead of collapsing to the + /// absent-axis default of 1. + pub fn init_ep(ep_size: usize, n_layers: usize) -> HipResult { + if ep_size == 0 { + return Err(HipError::new(0, "init_ep: ep_size must be >= 1")); + } + if n_layers == 0 { + return Err(HipError::new(0, "init_ep: n_layers must be >= 1")); + } + let device_ids = resolve_device_ids(ep_size)?; + let devices = construct_devices(&device_ids)?; + preflight_vram_with_opts(&devices, /*check_vram_delta=*/ true)?; + let band_starts = tp_band_starts(ep_size, n_layers); + + // PP=1 EP topology: every device runs every layer. Encode the layer + // map exactly as init_tp does so PP helpers stay well-defined, + // while the EP forward path dispatches every layer on every rank. + Ok(Self { + rccl_comms: None, + devices, + layer_to_device: vec![0u8; n_layers], + band_starts, + mesh: DeviceMesh::rect(&[(DimKind::Ep, ep_size)]).expect("ep mesh cannot overflow"), + peer_access_enabled: false, + output_device: 0, + givens_cos_per_dev: Vec::new(), + givens_sin_per_dev: Vec::new(), + peer_ar_tmp: Vec::new(), + peer_ar_tmp_bytes: 0, + host_ar_tmp: Vec::new(), + active_peer_lease: None, + peer_lease_buffers: Vec::new(), + peer_lease_next_id: 0, + peer_lease_quarantined: false, + rank_barrier_events: Vec::new(), + tp_graph_signals: Vec::new(), + tp_graph_barrier_count: 0, + tp_graph_capture_epoch: 0, + }) + } + /// Query whether every directed device pair supports peer access. /// Does not enable peer access or mutate peer state. Use when partial /// activation is unsafe (ROCm does not map later allocations). @@ -2138,4 +2187,21 @@ mod tests { assert_eq!(single.n_devices(), 1); assert!(single.axes().is_empty()); } + + #[test] + fn device_mesh_ep_group_and_tp_absent() { + // EP=4: the mesh `Gpus::init_ep` records. The Ep axis groups all + // four devices; the absent Tp axis reads back as 1 (the `size_of` + // default), so an EP load is never mistaken for TP. + let ep = DeviceMesh::rect(&[(DimKind::Ep, 4)]).unwrap(); + assert_eq!(ep.group_along(DimKind::Ep, &[0]).unwrap(), vec![0, 1, 2, 3]); + assert_eq!(ep.n_devices(), 4); + assert_eq!(ep.size_of(DimKind::Ep), 4); + assert_eq!(ep.size_of(DimKind::Tp), 1); + + // Symmetric check on the mesh `Gpus::init_tp` records. + let tp = DeviceMesh::rect(&[(DimKind::Tp, 4)]).unwrap(); + assert_eq!(tp.size_of(DimKind::Tp), 4); + assert_eq!(tp.size_of(DimKind::Ep), 1); + } } diff --git a/docs/CLI.md b/docs/CLI.md index 471fa9e1c..276351e57 100644 --- a/docs/CLI.md +++ b/docs/CLI.md @@ -121,7 +121,7 @@ Supported CLI formats include `mq4`, `mq6`, `q8`/`q8f16`, `hf4`/`hf6` and hfq al | Command | Purpose | |---|---| -| `hipfire bench [opts] [prompt]` | Prefill/decode timing. `--runs N` (default 5), `--json`, `--exp` (RDNA2 variant sweep). | +| `hipfire bench [opts] [prompt]` | Prefill/decode timing. `--runs N` (default 5), `--json`, `--exp` (RDNA2 variant sweep). `--prompt-file PATH` reads the prompt verbatim; JSON records `prompt_tokens`/`prompt_md5`/`prompt_chars`/`warnings` (short prompts warn that `prefill_tok_s` is launch overhead). | | `hipfire bench --matrix ...` | Synthetic PP/context/TG matrix (`--pp`, `--ctx`, `--tg`, `--sustained-tg`, `--sustained-ctx`, `--warmups`, `--kv-mode`, `--redline`). | | `hipfire profile [model] [--kernel substr] [--json]` | Live daemon roofline and compiled-kernel VGPR/SGPR/LDS/occupancy report. Use `hipfire-atlas` for measured ISA-fit and workload analysis. | | `hipfire diag` | Static device/runtime checks plus a live HIP arch, version, and VRAM probe when the daemon is available. | diff --git a/docs/MODELS.md b/docs/MODELS.md index b8efd4f1e..b46d3bc46 100644 --- a/docs/MODELS.md +++ b/docs/MODELS.md @@ -113,7 +113,7 @@ MQ2V2 is not registered. Explicit `qwen3.8:27b-mq4` aliases to `qwen3.8:27b`. Le | `qwen3.8:27b-draft-mq6` | `qwen38-27b-dflash-mq6.hfq` | 1.66 | 16 | `qwen3.8:27b*` (same-bit alt) | | `muse-glimmer:draft` | `muse-glimmer-30b-dflash.mq4` | 1.36 | 26 | `muse-glimmer` / `muse-glimmer:fast` | -Draft **loading** is controlled by `dflash_mode` / `speculation` / `HIPFIRE_DFLASH_DRAFT` ([`CONFIG.md`](CONFIG.md), [`env-vars.md`](env-vars.md)). Default `dflash_mode` is **off**. Filename auto-match may wire a sibling draft when present; that is discovery, not an admission that DFlash wins on every prompt. +Draft **loading** is registry-driven: `hipfire pull ` fetches the draft sidecar alongside the target, and `dflash_mode` / `speculation` ([`CONFIG.md`](CONFIG.md), [`env-vars.md`](env-vars.md)) decides what happens next — `auto` uses the sidecar when present (AR otherwise), `on` requires it (load fails without it), `off` never drafts. Default `dflash_mode` is **off**. `developer.dflash_draft` / `run --model-draft` override the sidecar. Watch for `DFlash draft loaded:` in load output. ### Qwen3 (non-3.5) dense HF4 diff --git a/docs/env-vars.md b/docs/env-vars.md index c2da45af4..fdfad68a1 100644 --- a/docs/env-vars.md +++ b/docs/env-vars.md @@ -105,7 +105,7 @@ Values and defaults below match `hipfire-config`, the native CLI, and/or `Runtim | Variable | Default / sense | Notes | |---|---|---| | `HIPFIRE_SPECULATION` | `off`/`auto`/`ngram`/`dflash`/`mtp`/`dspark` | Canonical selector | -| `HIPFIRE_DFLASH_DRAFT` | retired engine read | Still appears in legacy gate scripts; product draft discovery uses typed speculation/load policy and registry/filename matching. | +| `HIPFIRE_DFLASH_DRAFT` | explicit draft path (overrides the registry sidecar); empty opts out | Legacy `developer.dflash_draft` read; still appears in legacy gate scripts. | | `HIPFIRE_DFLASH_CTX_CAP` | **8192**; `0` restores uncapped legacy behavior | Caps draft-side context storage; over-cap requests fall back to AR | | `HIPFIRE_DFLASH_WINDOW` | **0 / unset** (legacy), unless declared by draft metadata | Enables bounded draft SWA; refused with CASK eviction | | `HIPFIRE_DFLASH_MODE` | RuntimeConfig default **`off`** | Distinct from config `dflash_mode` apply path — product CLI also uses load params | diff --git a/docs/quant-formats/mq4-v2.md b/docs/quant-formats/mq4-v2.md index 7c5d25fbb..3943cc21e 100644 --- a/docs/quant-formats/mq4-v2.md +++ b/docs/quant-formats/mq4-v2.md @@ -241,8 +241,11 @@ Measured across the 11 ported kernels, three distinct forms occur: | `kt < 8` | 8 | WMMA GEMM main + BT bodies | nibbles at `gp + 8 + kt*8 + k_grp*4` with `k_grp = tid>>4 ∈ {0,1}`; `kt*8 + 4 < 64` ⇔ `kt < 8` for both `k_grp` | | `quarter_in_group < 2` | 2 | WMMA `ldsstage` bodies | nibbles at `gp + 8 + quarter*32 + {0,8,16,24} + k_grp*4`; max offset at `quarter=1` is 60 < 64 | -In the WMMA main path the loop steps `kt += 4` over tiles `kt..kt+3`, so all four tiles in one -body share a half (kt=0,4 → half 0; kt=8,12 → half 1) and one select per body suffices. +The step varies by variant — gfx11 main bodies step `kt += 2` (K2 unroll; both tiles in a +pair share a half), the gfx11 residual body steps `kt++` (per-tile select), gfx12 bodies +step `kt += 4` over tiles `kt..kt+3` (all four tiles in one body share a half: kt=0,4 → +half 0; kt=8,12 → half 1) — but every body selects with `kt < 8`, so one select per +body suffices in all three forms. A wrong predicate here **compiles, runs, and silently applies the wrong scale to half of every tensor**. It is the single highest-risk detail in the port. For any kernel not in the 11, redo @@ -418,7 +421,7 @@ relatively less of the damage. "Codebooks are for the sub-4-bit tier" is defensi | **qt=45** batched prefill WMMA GEMM + batched lm_head GEMM | **gfx12-only** (`HasWmmaGfx12`); **not** admitted on gfx11 | | Unsupported batched prefill | **per-token decode fallback** (does not dispatch a foreign-arch WMMA kernel) | | FusedQkv / FusedGateUp decode registrations | remain **cross-arch** (not narrowed to gfx12) | -| Prefill LA admission (qt=44) | gfx1100/1101/1102/1150/1151 + gfx1200/1201; gfx11 opt-out `HIPFIRE_MQV2_GFX11_WMMA=0` | +| Prefill LA admission (qt=44) | **qwen35 only**: gfx1100/1101/1102/1150/1151 + gfx1200/1201; gfx11 opt-out `HIPFIRE_MQV2_GFX11_WMMA=0`. Plain-Llama dense (`llama::is_batchable_la`) **refuses V2 on every arch** (per-token decode) until `forward_prefill_chunk` grows V2 arms | | Exact parity examples | `mq4v2_parity`, `mq4v2_gemm_parity`, `mq4v2_fused_parity`, `mq4v2_residual_parity`, `mq4c_parity`; BT screens `test_mq4v2_*_bt_gfx{1100,1151,1201}.rs` | | Qwen3.8 fixture-bound KLD | qt=44 `ctl` WT2 0.039033 / v6 0.544517; `ctl2` WT2 0.032495; `attn` WT2 0.025437 (§ 5) | | gfx1010 | scalar fused decode TUs compile; batched prefill falls back as above | @@ -429,10 +432,10 @@ relatively less of the damage. "Codebooks are for the sub-4-bit tier" is defensi ### Not claimed / out of scope - **wave64** half-split (§ 4) — not verified; wave64 remains unsupported for these formats. -- **MoE** paths (`gemv_hfq4g256_moe_*`, `gemm_*_moe_grouped_*`) — out of scope / fail-closed for V2 product tiers. +- **MoE** paths (`gemv_mq4g256v2_moe_*`, `gemm_mq4g256v2_moe_grouped_*`) — production-wired for qt=44 (decode + prefill, gfx11 + gfx12; loader → `MoeResolution.routed_indexable_mq4v2` at `families/moe.rs:244` → `pipeline/mod.rs:1244-1262`). - **qt=45 on gfx11** — no gfx11 WMMA sibling; do not promote. - **gfx1030 default-R decision** (§ 6) — still open if/when gfx1030 ships these dtypes; not a dense-WMMA blocker. -- Research-only surfaces: `muse_*`, dp4a / cpol / `ldscoop` / `ldsx` / `.v1`–`.v5` / `XBATCH` single-row path. (gfx11 base/BT WMMA for **qt=44** is production, not research-only.) +- Research-only surfaces: `muse_*`, dp4a / cpol / `ldscoop` / `ldsx` / `.v1`–`.v5`. (The `XBATCH` single-row path IS ported for qt=44 — `gemv_mq4g256v2.hip:295-361` — and gfx11 base/BT WMMA for **qt=44** is production, not research-only.) ### Port surface that landed for dense HasWmma (qt=44) / dense gfx12 (qt=45) diff --git a/kernels/src/gemm_mq4g256v2_residual_mmq.hip b/kernels/src/gemm_mq4g256v2_residual_mmq.hip index 595eee5ff..43b3fcee4 100644 --- a/kernels/src/gemm_mq4g256v2_residual_mmq.hip +++ b/kernels/src/gemm_mq4g256v2_residual_mmq.hip @@ -7,7 +7,12 @@ #include #include -// Experimental HFQ4-G256 MMQ residual GEMM for RDNA3/RDNA3.5. +// Production MQ4G256V2 (qt=44) MMQ residual GEMM for gfx1100/gfx1151. +// +// Selected as the fast path when `batch >= 128 && batch % 128 == 0` +// (see `crates/rdna-compute/src/gemm.rs:17860-17909` and the qkv/qkvza/ +// gate_up/residual call sites); otherwise the WMMA residual path runs. +// Applies the scale per 128-K half (`kt < 8` select), not a per-256 V1-ism. // // This mirrors the important parts of llama.cpp's AMD MMQ path: // - pre-quantize the activation matrix into block_q8_1_mmq DS4 layout diff --git a/kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip b/kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip index 5959ae87f..8ec6bfa8d 100644 --- a/kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip +++ b/kernels/src/gemm_qkv_mq4g256v2_wmma.gfx12.hip @@ -12,17 +12,10 @@ // (which targets gfx11 / RDNA3). Compile this file with // hipcc --offload-arch=gfx1200 (or gfx1201, ...). // -// SCAFFOLD STATUS (2026-04-26): -// This file is NOT yet wired into dispatch.rs. It compiles for gfx12 -// targets and is intended as the canonical pattern reference for the -// five remaining gfx12 WMMA ports (qkvza-hfq4, gate_up-hfq4, plus -// the three hfq6 variants). Runtime correctness must be validated on -// real RDNA4 hardware (9070 XT / R9700) via test_kernels before any -// dispatch.rs site is allowed to route here. The C-output mapping -// below is a HYPOTHESIS derived from CK trait math and has NOT been -// confirmed against silicon — see the channel-test note at the -// write-back loop. See `.skills/hipfire-arch-port/` for the port -// workflow. +// STATUS: production gfx12 WMMA path, wired into dispatch. The C-output +// mapping below was validated 2026-04-27 in this file's channel-test on +// R9700 — the same map its `gemm_gate_up_mq4g256v2_wmma.gfx12.hip` sibling +// documents. // // Differences from the gfx11 kernel: // 1. WMMA builtin: __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12 @@ -34,11 +27,12 @@ // lane (tid >> 4) = 0 -> carries K = [0..7] of the 16-K tile // lane (tid >> 4) = 1 -> carries K = [8..15] of the 16-K tile // Each lane therefore loads HALF as much per WMMA tile as gfx11. -// 4. C-output mapping (HYPOTHESIS — see below): -// gfx11 (validated, commit b7ac66a): -// acc[j] = C[2*j + (tid>>4)][tid & 15] -// gfx12 (UNVERIFIED, derived from CK kCM0/kCM1PerLane swap): -// acc[j] = C[8*(tid>>4) + j][tid & 15] +// 4. C-output mapping (validated 2026-04-27 in the QKV scaffold's +// channel-test on R9700): +// gfx11: acc[j] = C[2*j + (tid>>4)][tid & 15] (rows interleaved) +// gfx12: acc[j] = C[8*(tid>>4) + j][tid & 15] (rows contiguous — +// group 0: rows 0..7, +// group 1: rows 8..15) // // Reference (ROCm 7.x): // /opt/rocm/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_wmma_impl_base_traits.hpp @@ -165,10 +159,8 @@ extern "C" __global__ void gemm_qkv_mq4g256v2_wmma_gfx12( } // --- Output --- - // - // gfx12 wave32 WMMA C-mapping HYPOTHESIS — DO NOT TRUST WITHOUT - // CHANNEL-TEST ON REAL HARDWARE. Derived from the CK trait swap - // (warp_gemm_attribute_wmma_impl_base_traits.hpp): + // gfx12 wave32 WMMA C-mapping (validated 2026-04-27 in this file's + // channel-test on R9700; same map the gate_up gfx12 sibling documents): // // gfx11: kCMLane=2, kCM0PerLane=8, kCM1PerLane=1 // -> acc[j] = C[2*j + (tid>>4)][tid & 15] (rows interleaved) @@ -176,12 +168,6 @@ extern "C" __global__ void gemm_qkv_mq4g256v2_wmma_gfx12( // -> acc[j] = C[8*(tid>>4) + j][tid & 15] (rows contiguous) // // i.e. lane group 0 holds output rows 0..7, lane group 1 rows 8..15. - // The gfx11 mapping was silently wrong for ~6 weeks before being fixed - // in commit b7ac66a — assume the same risk here. Validation recipe - // (from that fix): add `if (blockIdx.x == 0 && blockIdx.y == 0) - // printf("tid=%d j=%d row=%d col=%d acc=%f\\n", tid, j, out_row, - // out_col, acc[j]);` and compare against a CPU reference for a small - // golden case. Adjust the (j, k_grp) -> out_row formula until it matches. const int out_col = batch_start + m_lane; // batch index if (out_col < N) { #pragma unroll diff --git a/registry/models.json b/registry/models.json index d18e5332e..e7a98751c 100644 --- a/registry/models.json +++ b/registry/models.json @@ -102,6 +102,9 @@ "presence_penalty": 1.5, "repeat_penalty": 1.0 } + }, + "dflash": { + "file": "qwen35-9b-dflash-mq4.hfq" } }, "qwen3.5:27b": { @@ -144,6 +147,9 @@ "presence_penalty": 1.5, "repeat_penalty": 1.0 } + }, + "dflash": { + "file": "qwen35-27b-dflash-mq4.hfq" } }, "qwen3.5:35b-a3b": { @@ -697,6 +703,9 @@ "desc": "44 tok/s AR / 185 tok/s w/ draft on code", "triattn": { "file": "qwen3.6-27b.mq4.triattn.blended_v3.bin" + }, + "dflash": { + "file": "qwen36-27b-dflash-mq4.hfq" } }, "qwen3.8:27b-mq3-xt": { @@ -725,6 +734,9 @@ "repeat_penalty": 1.0, "reasoning_effort": "xhigh" } + }, + "dflash": { + "file": "qwen38-27b-dflash-mq3.hfq" } }, "qwen3.8:27b-mq3": { @@ -753,6 +765,9 @@ "repeat_penalty": 1.0, "reasoning_effort": "xhigh" } + }, + "dflash": { + "file": "qwen38-27b-dflash-mq3.hfq" } }, "qwen3.8:27b-mq3-pro": { @@ -781,6 +796,9 @@ "repeat_penalty": 1.0, "reasoning_effort": "xhigh" } + }, + "dflash": { + "file": "qwen38-27b-dflash-mq3.hfq" } }, "qwen3.8:27b-mq4-xt": { @@ -809,6 +827,9 @@ "repeat_penalty": 1.0, "reasoning_effort": "xhigh" } + }, + "dflash": { + "file": "qwen38-27b-dflash-mq4.hfq" } }, "qwen3.8:27b": { @@ -837,6 +858,9 @@ "repeat_penalty": 1.0, "reasoning_effort": "xhigh" } + }, + "dflash": { + "file": "qwen38-27b-dflash-mq4.hfq" } }, "qwen3.8:27b-mq4-pro": { @@ -865,6 +889,9 @@ "repeat_penalty": 1.0, "reasoning_effort": "xhigh" } + }, + "dflash": { + "file": "qwen38-27b-dflash-mq4.hfq" } }, "qwen3.8:27b-mq5-xt": { @@ -1060,7 +1087,10 @@ "min_vram_gb": 8.8, "desc": "MQ6, higher quality", "default_kv_mode": "q8", - "quant_recipe": "v3-awq-f1" + "quant_recipe": "v3-awq-f1", + "dflash": { + "file": "qwen35-9b-dflash-mq4.hfq" + } }, "qwen3.5:27b-mq6": { "repo": "hipfire-models/qwen3.5-27b", @@ -1076,14 +1106,20 @@ "min_vram_gb": 6.1, "desc": "MQ3 alpha (3.25 bpw, gfx11/gfx12). Smaller than MQ4, comparable decode. Quality eval pending — see issue #113. Sub-9B MQ3 not shipped — see #114.", "default_kv_mode": "q8", - "quant_recipe": "v3-awq-f1" + "quant_recipe": "v3-awq-f1", + "dflash": { + "file": "qwen35-9b-dflash-mq4.hfq" + } }, "qwen3.5:27b-mq3": { "repo": "hipfire-models/qwen3.5-27b", "file": "qwen3.5-27b.mq3", "size_gb": 10.7, "min_vram_gb": 12, - "desc": "MQ3 alpha — fits 128K asym3 ctx on 24 GB (MQ4 OOMs at ~98K). gfx11/gfx12 only." + "desc": "MQ3 alpha — fits 128K asym3 ctx on 24 GB (MQ4 OOMs at ~98K). gfx11/gfx12 only.", + "dflash": { + "file": "qwen35-27b-dflash-mq3.hfq" + } }, "qwen3.6:27b-mq3": { "recommended_settings": { @@ -1096,7 +1132,10 @@ "file": "qwen3.6-27b.mq3", "size_gb": 10.7, "min_vram_gb": 12, - "desc": "MQ3 alpha — fits 128K asym3 ctx on 24 GB. Pairs well with mq4 DFlash draft (126 tok/s τ=7.0)." + "desc": "MQ3 alpha — fits 128K asym3 ctx on 24 GB. Pairs well with mq4 DFlash draft (126 tok/s τ=7.0).", + "dflash": { + "file": "qwen36-27b-dflash-mq3.hfq" + } }, "qwen3.5:27b-draft-mq3": { "repo": "hipfire-models/qwen3.5-27b", diff --git a/registry/v1.json b/registry/v1.json index 65aa9982a..23cca9525 100644 --- a/registry/v1.json +++ b/registry/v1.json @@ -113,6 +113,11 @@ "repeat_penalty": 1.0 } }, + "dflash": { + "file": "qwen35-9b-dflash-mq4.hfq", + "sha256": "97457f684b8f0b195e631025a63da0a92286d1c77751e58caf0cd470457a6b25", + "size_bytes": 557265920 + }, "sha256": "ba83acf5bfd5d4e334b0afc26d779734e31623bb7f74e807c3581dfecb3128ad", "size_bytes": 5313750016, "arch_id": 5, @@ -159,6 +164,11 @@ "repeat_penalty": 1.0 } }, + "dflash": { + "file": "qwen35-27b-dflash-mq4.hfq", + "sha256": "3d428b97c1911a9ad815cc52fbee080306852c1dafad6b1b17bb70bd68010301", + "size_bytes": 919401472 + }, "sha256": "ea615949ddf6a180eee03ff6fde39f7e51148f153b1b05f82258b9953088576e", "size_bytes": 14984158208, "arch_id": 5, @@ -792,6 +802,11 @@ "sha256": "d6cb8026841830cfeb82d2709453aa753f65b5596bfb9cc9c085c808fda6ad22", "size_bytes": 2359324 }, + "dflash": { + "file": "qwen36-27b-dflash-mq4.hfq", + "sha256": "bd8c4f07ae80fe1385bf2606af9a7ba0daa18ca8daec50916f2a489054c44e70", + "size_bytes": 919401472 + }, "sha256": "86a5f80fd29d545abb1093dead242725ced6d68b8607c6d566d897b1a82442dc", "size_bytes": 14984158208, "arch_id": 5, @@ -824,6 +839,11 @@ "reasoning_effort": "xhigh" } }, + "dflash": { + "file": "qwen38-27b-dflash-mq3.hfq", + "sha256": "dd335fb2634c371d120077c2fe754643c91b27f32d0c32668278fe4b8f1ca319", + "size_bytes": 984978432 + }, "sha256": "3e04fc8db80bda557b965ec60ac876cf2500fced7f340624f3fcbeae134af5c5", "size_bytes": 11777616896, "arch_id": 5, @@ -856,6 +876,11 @@ "reasoning_effort": "xhigh" } }, + "dflash": { + "file": "qwen38-27b-dflash-mq3.hfq", + "sha256": "dd335fb2634c371d120077c2fe754643c91b27f32d0c32668278fe4b8f1ca319", + "size_bytes": 984978432 + }, "sha256": "09c3544690aceca29e1822d79adab6ffcc8fd9e4b58359fe8dfb185ef49811c9", "size_bytes": 12618796032, "arch_id": 5, @@ -888,6 +913,11 @@ "reasoning_effort": "xhigh" } }, + "dflash": { + "file": "qwen38-27b-dflash-mq3.hfq", + "sha256": "dd335fb2634c371d120077c2fe754643c91b27f32d0c32668278fe4b8f1ca319", + "size_bytes": 984978432 + }, "sha256": "394c50966bf4f68172df8eb34cd7ded8f9d0576c9ef24ea6ba639a88c184f795", "size_bytes": 13184433152, "arch_id": 5, @@ -920,6 +950,11 @@ "reasoning_effort": "xhigh" } }, + "dflash": { + "file": "qwen38-27b-dflash-mq4.hfq", + "sha256": "d0a74a232a0e2166d889f823e91e0fbf778d21dd9668d7de055cdecb065401bc", + "size_bytes": 1209603072 + }, "sha256": "9f91556f7e0431a077d03756a7102d0154108757289e6e5fe9a2d204c0c9eeb7", "size_bytes": 14980361216, "arch_id": 5, @@ -952,6 +987,11 @@ "reasoning_effort": "xhigh" } }, + "dflash": { + "file": "qwen38-27b-dflash-mq4.hfq", + "sha256": "d0a74a232a0e2166d889f823e91e0fbf778d21dd9668d7de055cdecb065401bc", + "size_bytes": 1209603072 + }, "sha256": "5bb556a6cc84035234995c017c9791aa3951ad1eae4cf8c8172b0eaef399e507", "size_bytes": 15662615552, "arch_id": 5, @@ -984,6 +1024,11 @@ "reasoning_effort": "xhigh" } }, + "dflash": { + "file": "qwen38-27b-dflash-mq4.hfq", + "sha256": "d0a74a232a0e2166d889f823e91e0fbf778d21dd9668d7de055cdecb065401bc", + "size_bytes": 1209603072 + }, "sha256": "e6f2ac87042b9e314c323f00bc499a6304cd5624021ae501ce90b12a3a7ea3fa", "size_bytes": 16464182272, "arch_id": 5, @@ -1215,6 +1260,11 @@ "desc": "MQ6, higher quality", "default_kv_mode": "q8", "quant_recipe": "v3-awq-f1", + "dflash": { + "file": "qwen35-9b-dflash-mq4.hfq", + "sha256": "97457f684b8f0b195e631025a63da0a92286d1c77751e58caf0cd470457a6b25", + "size_bytes": 557265920 + }, "sha256": "69b0e3b2be99a7fcab17f82bae2a2f1342ac32ee96d421347726815a69e78ce4", "size_bytes": 7296132096, "arch_id": 5, @@ -1239,6 +1289,11 @@ "desc": "MQ3 alpha (3.25 bpw, gfx11/gfx12). Smaller than MQ4, comparable decode. Quality eval pending \u2014 see issue #113. Sub-9B MQ3 not shipped \u2014 see #114.", "default_kv_mode": "q8", "quant_recipe": "v3-awq-f1", + "dflash": { + "file": "qwen35-9b-dflash-mq4.hfq", + "sha256": "97457f684b8f0b195e631025a63da0a92286d1c77751e58caf0cd470457a6b25", + "size_bytes": 557265920 + }, "sha256": "c379dbbc90d7faf5e7281f01310b4e3f3e76587a6e951a0c7b6a809eeae5550b", "size_bytes": 4569785344, "arch_id": 5, @@ -1250,6 +1305,11 @@ "size_gb": 10.7, "min_vram_gb": 12, "desc": "MQ3 alpha \u2014 fits 128K asym3 ctx on 24 GB (MQ4 OOMs at ~98K). gfx11/gfx12 only.", + "dflash": { + "file": "qwen35-27b-dflash-mq3.hfq", + "sha256": "b3526f2f8ffce5a2483f6d74fd1e78af2b8eebe9249e5b55177ed0e0c58f53fc", + "size_bytes": 703132672 + }, "sha256": "58fdaf54ec3c4be9372bd44b4754e198e7ff6d7129c6abdaec76c1ac70fb092b", "size_bytes": 11784735744, "arch_id": 5, @@ -1267,6 +1327,11 @@ "size_gb": 10.7, "min_vram_gb": 12, "desc": "MQ3 alpha \u2014 fits 128K asym3 ctx on 24 GB. Pairs well with mq4 DFlash draft (126 tok/s \u03c4=7.0).", + "dflash": { + "file": "qwen36-27b-dflash-mq3.hfq", + "sha256": "f5fc4bef4c15229f940f8ead6c04723b04308321c4a9584c3318d364cadf1728", + "size_bytes": 703132672 + }, "sha256": "6b650a7d6cb4e2447dd67da7a87cb5e57956e98eff676a498a7586e3b3cc9298", "size_bytes": 11776322560, "arch_id": 5, diff --git a/scripts/registry_gen.py b/scripts/registry_gen.py index eef423544..069f795a3 100644 --- a/scripts/registry_gen.py +++ b/scripts/registry_gen.py @@ -12,7 +12,7 @@ purely additive: top-level : schema_version, generated_at per-entry : sha256 (HF LFS oid), size_bytes, arch_id, quant - sidecars : triattn/mtp gain sha256/size_bytes next to their `file` + sidecars : triattn/mtp/dflash gain sha256/size_bytes next to their `file` Fail-closed: ANY problem — repo unreachable, file missing from the repo tree, file not LFS (no sha256), size_bytes disagreeing with curated @@ -360,7 +360,7 @@ def is_strict_superset(old: object, new: object, path: str, errors: list[str]) - def annotate_sidecar( sidecar: dict, tree: dict[str, dict], tag: str, kind: str, errors: list[str] ) -> dict: - """triattn/mtp sub-object: require existence, add sha256/size_bytes if LFS.""" + """triattn/mtp/dflash sub-object: require existence, add sha256/size_bytes if LFS.""" out = dict(sidecar) fname = sidecar.get("file", "") item = tree.get(fname) @@ -482,7 +482,7 @@ def build_registry(curated: dict, token: str | None) -> tuple[dict | None, list[ f"HF {size_bytes / 1e9:.2f} GB ({drift:.0%} drift); " f"update registry/models.json" ) - for kind in ("triattn", "mtp"): + for kind in ("triattn", "mtp", "dflash"): if isinstance(entry.get(kind), dict): new_entry[kind] = annotate_sidecar(entry[kind], tree, tag, kind, errors) # repo probe already failed → error recorded above; entry still gets