v0.2.1 — Dispatch unification

v0.2.1 — Dispatch unification (#397)

The centralized kernel-dispatch program lands: every GEMV/GEMM/attention/
MoE/rotation/fused projection across qwen35, llama, qwen2, dots-ocr,
deepseek4, minimax and lfm2moe now resolves through typed kernel families
with per-(arch x dtype) tables — a new dense quant is a table entry plus a
kernel file, no model code. The lowered forward-as-pipeline decode is
default-on for all seven arches (qwen2/dots-ocr byte-parity validated on
gfx1100 + gfx1201). KvTierPlan unifies KV-write/flash-attend resolution
with adaptive-KV exemptions; GPU-free coverage gates assert dispatch plans
for the entire fleet including RDNA4 rows.

Also in this release: gfx12 F16 GEMV fix (broken gemm_f16_tiled fallback
rewrote ds4 DSA compressor garbage on RDNA4 — restored byte-parity with
RDNA3.5 and ds4 EP tool-calling); gemm_f16_tiled rewritten (correct + up
to 13.8x faster); jinja chat templates default-on; DSML render/parse on
the EP serve path; q8 error-feedback DeltaNet state default; master line
fully merged back (hunt-3 fixes, FP32/Q4 DeltaNet spec-decode, MQ6 DFlash,
F32 oracle passthrough quantizer mode).

v0.2.0 — DeepSeek V4 Flash + tokenizer diagnostics

v0.2.0 — DeepSeek V4 Flash + tokenizer diagnostics

DeepSeek V4 Flash is now a first-class hipfire architecture (arch_id=9).
The new hipfire-arch-deepseek4 crate wires the production
deepseek-v4-flash.mq2lloyd path through the canonical daemon and CLI
surface: hipfire-quantize → hipfire serve → hipfire run, with no Python
in the hot path.

DeepSeek V4 Flash

• New architecture crate: config parsing, weight loading, runtime state,
  prefill/decode, and MTP speculative decode live under
  crates/hipfire-arch-deepseek4/.
• New DeepSeek V4 kernel surface: sliding-window attention, compressed-KV
  indexer, Hyper-Connections, compressor/indexer projections, MQ2-Lloyd MoE
  GEMVs, tail-only YaRN RoPE, and glue kernels.
• New quantizer formats for DeepSeek V4 source / Q8 / Q8+MTP packaging, plus
  hfq_split for moving mtp.0.* tensors into an optional sidecar so normal
  decode does not upload the extra MTP layer.
• Daemon support for arch_id=9 includes batched prefill, deterministic
  routing defaults, plain decode, and MTP speculative decode controlled by
  HIPFIRE_DEEPSEEK4_SPEC_DECODE, HIPFIRE_DEEPSEEK4_SPEC_K, and
  HIPFIRE_DEEPSEEK4_MTP_ADDON.
• Tool-call output works on both plain decode and the DeepSeek MTP path: the
  daemon emits tool_calls events with finish_reason: "tool_calls" instead
  of leaking raw DSML/tool-call text.

Validated on gfx1151 / Radeon 8060S with HIP_VISIBLE_DEVICES=1:

• cargo test -p hipfire-arch-deepseek4 --lib
• cargo check -p hipfire-arch-deepseek4 --examples
• cargo check --workspace --examples
• ./scripts/coherence-gate.sh --full
• ./scripts/coherence-gate-deepseek4-mtp.sh --full

Tokenizer interned symbols + loud OOV at construction

Tokenizer::from_* constructors now return Result<Self, TokenizerError>
instead of Option<Self>. Inconsistent vocab/merges pairs (e.g. truncated
quantizer output, vocab missing a byte char, merges referencing absent
symbols) are now rejected loudly at load time with a specific error variant
instead of silently producing a Tokenizer whose encode_gpt2_bpe would
emit id 0 for OOV symbols downstream (#203).

The internal merge representation is now token-id-keyed end-to-end. The GPT-2
BPE encoder operates on Vec<u32> (was Vec<String>); merge-pair lookups
use HashMap<(u32, u32), u32> (was HashMap<(String, String), usize>).
Heap-loop String clones are eliminated. For a Qwen3-class vocab this saves
roughly 13 MB per loaded tokenizer.

Public API change

Old                                          → New
Tokenizer::from_gguf(...) -> Option<Self>    → Result<Self, TokenizerError>
Tokenizer::from_hf_json(...) -> Option<Self>  → Result<Self, TokenizerError>
Tokenizer::from_hfq_metadata(...)            → Result<Self, TokenizerError>
Tokenizer::from_gguf_meta_json(...)          → Result<Self, TokenizerError>
Speculative::load_tokenizer(&self) -> Option<Tokenizer>
                                             → Result<Tokenizer, TokenizerError>

New public types in hipfire_runtime::tokenizer:

• TokenizerError — variants: MetadataMissing { field }, MalformedJson,
  MissingByteSymbol { byte, char }, MissingMergeOperand { rank, left, right, missing_side }, MissingMergeResult { rank, expected }. Implements
  Display, std::error::Error, From<serde_json::Error>.
• Side — Left | Right, used by MissingMergeOperand.

Migration guide for contributors

Old caller pattern                           → New caller pattern
.expect("...")                                (unchanged — works on Result)
.unwrap()                                     (unchanged — works on Result)
.ok_or(_)?  / .ok_or_else(|| _)?              .map_err(|e| ...: {e})?
                                              or .map_err(|_| _)?
if let Some(t) = ...from_*(...)               if let Ok(t) = ...from_*(...)
.unwrap_or_else(|| ...)                       .unwrap_or_else(|_| ...)
                                              (closure now receives error)

Why

Pre-existing Option-returning constructors made every failure mode look
identical from the outside. A user whose model failed to load got None
with no information about why. With many possible failure modes (corrupt
JSON, missing metadata, vocab/merges drift after re-quantization), this
made remote debugging painful. The new Result variants describe exactly
what's wrong (which byte, which merge rank, which symbol), so a single log
line is enough to diagnose.

The OOV consistency check at construction is the structural fix for #203.
The encoder no longer needs to silently fall back to id 0 in its final walk
because the constructor guarantees no OOV symbol can survive merges.

Caveats

• The SentencePiece encoder's single-char fallback at best_len == 0
  still silently drops missing chars. That's a separate failure mode that
  needs an encode_strict variant returning Result<Vec<u32>, EncodeError>;
  deferred to a follow-up PR.
• The max_token_chars cap on the SentencePiece greedy scan (bounding the
  pre-existing O(N²) tail on unmatchable inputs) is also deferred.

v0.1.20 — engine modularization

Major architectural refactor. The monolithic `engine` crate has split
into a `hipfire-runtime` crate plus per-arch crates. New `Architecture`
trait makes adding model archs a clean bring-up workflow. Performance
preserved within ±2% across all hosts; decode is bit-identical at 9B
on gfx1151 (Strix Halo).

Discord: https://discord.gg/F3BaywB8Rs for real-time chat,
contributor coordination, rebase help.

Highlights

• engine crate → hipfire-runtime (new name, same role + four new
  modules)
• Per-arch crates: `hipfire-arch-qwen35`, `hipfire-arch-qwen35-vl`,
  `hipfire-arch-llama`, `hipfire-arch-toy` (template for new arches)
• `Architecture` trait scaffold in
  `crates/hipfire-runtime/src/arch.rs` — bring-up spec for new model arches
• Generation guards extracted into reusable modules: `loop_guard`
  (n-gram detector), `sampler` (top-p / repeat-penalty / blocked-tokens),
  `prompt_frame` (ChatML), `eos_filter` (output stream)
• Bare examples gain `--guards on/off` flag — `infer_qwen35` /
  `infer_qwen3` can now opt into production-quality generation hygiene
  while preserving bare-engine probe semantics by default
• Contributor onboarding: CONTRIBUTING.md crate-topology section,
  CHANGELOG migration map, PR template, `hipfire-arch-toy` reference
  crate as a copy-paste starting point for new arches
• Rebase helper: `scripts/rebase-onto-modular.sh` +
  `.skills/rebase-onto-modular/SKILL.md` mechanically port pre-modular
  branches to the new topology

Breaking — for contributors with in-flight branches

If you have an open PR or local branch authored against pre-0.1.20
master, you'll see import / path errors after rebase. The new tooling
handles ~80% mechanically:

```bash
git checkout your-feature-branch
./scripts/rebase-onto-modular.sh
```

The script creates a backup tag, rebases onto master, applies the
path-rename + import-rewrite map, and reports any remaining manual
fixes. See issue #155
for the full migration guide and common-conflict table.

Quick old → new path map

Old	New
`crates/engine/src/qwen35.rs`	`crates/hipfire-arch-qwen35/src/qwen35.rs`
`crates/engine/src/qwen35_vl.rs`	`crates/hipfire-arch-qwen35-vl/src/qwen35_vl.rs`
`crates/engine/src/llama.rs`	`crates/hipfire-runtime/src/llama.rs` (facade)
`crates/engine/src/speculative.rs`	`crates/hipfire-arch-qwen35/src/speculative.rs`
`crates/engine/src/pflash.rs`	`crates/hipfire-arch-qwen35/src/pflash.rs`
`use engine::*`	`use hipfire_runtime::*` (or arch crate)
`cargo build -p engine`	`cargo build -p hipfire-runtime`

Full map in CHANGELOG.md.

What didn't change

• Kernel files (`kernels/src/*.hip`) — unchanged
• `rdna-compute/dispatch.rs` per-RDNA-arch routing — unchanged
• `hip-bridge` HIP/ROCm FFI — unchanged
• `hipfire-quantize` CLI — unchanged
• Daemon JSON-line API surface — unchanged
• All 45 `HIPFIRE_*` environment variables — preserved bit-for-bit
• Locked speed-gate baselines — perf parity verified on gfx1100 + gfx1151

Performance verification (canonical speed-gate config)

`HIPFIRE_KV_MODE=asym3 HIPFIRE_DPM_WARMUP_SECS=3 HIPFIRE_GRAPH=1` against
`qwen3.5-9b.mq4` on freshly purged + recompiled kernel cache:

Host	Metric	Master	Integration	Δ
gfx1100 (7900 XTX)	gen_tok_s	123.1	124.5	+1.1%
gfx1100 (7900 XTX)	prefill_tok_s	1809	1821	+0.6%
gfx1151 (Strix Halo)	gen_tok_s	45.5	45.5	0.0%
gfx1151 (Strix Halo)	bw_gib_s	225.1	225.3	+0.1%

Kernel hash sum (48 .hash files, gfx1151, fresh recompile): byte-identical
at `09bc045fe3db6aa81f87abce03cd1be6ec528e9666b08e457d2f7585979646e6`
across master and integration daemons.

What's still TODO post-0.1.20

• Transformer-extraction PR — pulls cross-arch primitives
  (`weight_gemv`, `KvCache`, dequant helpers, RoPE) out of `llama.rs`
  into `hipfire_runtime::transformer::*`. Unblocks the physical
  llama-arch crate split (currently a facade).
• Gemma branch forward-port — `gemma4` lives on its own branch with
  pre-modular daemon hooks. Forward-port happens after transformer
  extraction lands.
• `docs/architecture-ids.md` — arch_id registry to prevent
  collisions across community contributions.

Adding a new model architecture

`crates/hipfire-arch-toy/` is a minimal stub arch (~150 lines, 4 doc-
heavy files) demonstrating `Architecture` trait impl. Copy as starting
point for a new arch; see CONTRIBUTING.md "Crate topology" + decision
tree. RDNA GPU arch ports remain a separate concern handled by
`.skills/hipfire-arch-port/` (kernels live in `kernels/src/`,
dispatch in `rdna-compute`).

Issues / PRs

• Pinned migration guide: #155
• PRs in flight that may need rebase: #147 (UMA loader/MQ6),
  #148 (OpenAI streaming stats), #153 (weight-pager v0.2)

— @Kaden-Schutt

v0.1.9-alpha.1 — #111 tool-call attractor block

v0.1.9-alpha.1 (2026-05-02)

Patch release. Closes #111 — MQ4 single-token attractor on <tool_call>
that left agentic harnesses unable to dispatch tool calls on
qwen3.6:27b.mq4 (and any other MQ4-quant Qwen3+ model with structured-
output drift). Two complementary defenses ship together:

• Engine (daemon) — new GPU-side apply_unclosed_attractor_block
  scans the recent decode window for unclosed <tool_call> openers
  (opens − closes); when depth ≥ 2, writes a single 4-byte -INF
  to the logits buffer at the opener's token offset before the next
  gpu.sample_top_p. Same gate for <think> to head off the
  thinking-mode boundary corruption the same reporter saw. Cost is
  zero when not tripped, ~5 µs when tripped — no D2H, no kernel
  change. The (opens − closes) invariant means legitimate multi-
  tool turns never trip: a complete <tool_call>...</tool_call>
  decrements depth before the next opener arrives.

• CLI (parser) — parseToolCalls now strips any leading
  <tool_call>\s* repeats from a captured block before JSON parse.
  Defense-in-depth in case a nested opener does slip through (the
  engine block fires before the third opener, but the second still
  ships in the visible stream).

Verified on hardware (gfx1100 / 7900 XTX / ROCm 7.2 / qwen3.6:27b.mq4)
against the two prompts the reporter posted as still-broken after the
v0.1.9-alpha defensive parser ship:

Prompt 1: "what files are in this directory?"
→ tool_calls: [{ name: "bash", arguments: { command: "ls -la" } }]

Prompt 2: "write a file here named test.md with the text test inside"
→ tool_calls: [{ name: "write", arguments: { path: "test.md", content: "test" } }]

Both finish_reason: "tool_calls". No attractor loop, no nested
openers, no parser repair signal on stderr. The model emits clean spec
JSON now that the gate is in place.

Tests: 10 Rust unit tests (llama::tests) covering threshold edges,
window scope, complete-pair-allow, depth-saturate-at-zero. 14 Bun tests
(cli/parse_tool_calls.test.ts) including 4 new for nested-opener
strip. Coherence-gate green on 6/6 prompts including the existing
tool-call coverage test.

Still a stopgap on the symptom. The underlying root cause is MQ4
calibration drift on structured-output token positions; Path C
calibration retrain (#39) is the proper fix.

v0.1.9-alpha — MQ3 production-ready

v0.1.9-alpha (2026-05-02)

Headline: MQ3 is production-ready. The sub-4-bit Magnum Quant from
v0.1.8-alpha is now a full first-class citizen alongside MQ4 — K4-unrolled
decode GEMV, WMMA prefill family, DFlash cross-quant matrix, gfx12 port.
27B MQ3 fits 128K context in 24 GB where MQ4 OOMs at ~115K. Plus six
contributor PRs land in the same cycle, two arch bring-ups, and a sweep
of cache-lifecycle and parser hardening in response to user reports.

Highlights — MQ3 production push

• K4-unrolled MQ3 decode GEMV + fused residual (gfx1100). 9B MQ3
  decode 114 → 141 tok/s (+24%); 4B and 0.8B see proportional wins. Same
  pattern as the v0.1.8 K4 unroll on HFQ4: 4 weight reads + 4 X reads
  hoisted, 4 dequant + accumulate pairs in the body. Kernel matches MQ4
  decode within 2% on every size despite the 104 vs 136 B/group.
• WMMA prefill family for HFQ3 — gemm_qkvza_hfq3g256_wmma,
  gemm_qkv_hfq3g256_wmma, gemm_gate_up_hfq3g256_wmma,
  gemm_hfq3g256_residual_wmma. Closes the 17× prefill gap that gated
  ship: 9B MQ3 pp32 962 tok/s, pp128 1527 tok/s. Arch-gated to gfx11
  wave32 WMMA (gfx1100/1101/1102/1150/1151); gfx12 K4 variant landed in
  this same cycle.
• gfx12 (RDNA4) MQ3 WMMA port — full 4-kernel family ported to
  _w32_gfx12 builtin with K4 unroll + half8_t lane-split matching the
  v0.1.8 HFQ4 work. gfx1201 baseline + speed-baselines committed.
• DFlash + MQ3 cross-quant matrix. MQ3-target ↔ MQ3-draft, MQ3-target
  ↔ MQ4-draft, MQ4-target ↔ MQ3-draft all validated end-to-end on
  gfx1100. Refusal logic narrowed: MoE/A3B + MQ3 still refused (no MoE
  branched WMMA path); dense MQ3 ships. CLI auto-discovery prefers
  dirname(target) first then mq3↔mq4 cross-quant fallback dirs.
• 27B MQ3 context-fit data — fits 128K context in 24 GB on gfx1100
  with asym3 KV (10.7 GiB weights vs MQ4's 13.4 GiB). The 2.7 GiB
  saved on weights is the difference between fitting 128K and OOMing
  at 112K.

Highlights — contributor PRs

• PR #118 — Per-weight MMQ auto-dispatch (@fivetide). HFQ4 prefill
  routes to the MMQ i8-WMMA path automatically when batch_size ≥ 256
  and arch supports it (gfx1100/1101/1102/1103/1150/1151/1152). 9B
  pp512 +27% (1672 → 2122 tok/s). Default-on; opt out with
  HIPFIRE_MMQ=off. Tri-state config (off/on/auto) exposed in
  hipfire-tui.
• PR #117 — Windows serve -d parity (@fivetide). hipfire serve
  daemon mode wired through PowerShell on Windows; matches the Linux
  --detach UX. Includes compile-kernels.ps1 (PowerShell port of
  compile-kernels.sh), install.ps1 parity with daemon precompile,
  and hipcc.exe-first preference for paths with spaces.
• PR #103 — Raw-filename → registry-tag (@fivetide). hipfire run qwen3.5-9b.mq4 "..." now resolves the per-model overrides
  (max_think_tokens, temp, etc.) the same way as the registry-tag
  form. Prior behavior silently fell back to global defaults.
• PR #91 — gfx12 WMMA K4 K-tile unroll (@RobinVanCauter). 7
  .gfx12.hip kernels rewritten with kt += 4 inner loop;
  tests/speed-baselines/gfx1201.txt refreshed; closes #65. Includes
  bench-cold.sh for cold-process N-run distribution capture.
• PR #93 — gfx906 / Vega 20 / MI50 bring-up (@myaple). Family-141
  rev-detect splits gfx906 from gfx900 in redline::device; wave64
  dispatch list extended to gfx906; compile-kernels.sh skips
  WMMA/dot8 on gfx906; HSA_OVERRIDE + rocminfo fallbacks added. 196/196
  kernel compiles + 16/16 channel-tests on local MI50.
• PR #109 — MQ2 refuse + MQ3 advisory + sweep harness (mine, Codex
  follow-ups). MQ2 refused-by-default (severe quality cliff confirmed);
  MQ3 emits an advisory on sub-9B; sweep harness reproducibility per
  CLAUDE.md (committed prompts as files with md5 manifest, scratch off
  /tmp).

Engine

• Cache-invalidation lifecycle (Codex stop-time follow-ups). Three
  cross-cutting issues fixed:
  • Gpu::invalidate_weight_caches() clears mmq_screen_cache and
    drains fp16_shadow_cache on unload_model. Previous behavior left
    pointer-keyed cache hits on freed buffers — silent corruption on
    next model load if HIP reused the address.
  • Gpu::invalidate_graph_state() calls graph_destroy + verify_graph_destroy_all + replay_graph_destroy_all on unload.
    Captured hipGraphs over freed weight tensors would replay against
    garbage on the next forward_scratch_warmed_up call.
  • graph_destroy() resets ar_forward_warmed_up = false. Without
    this, the next forward_scratch would skip the warmup path and try
    to replay a destroyed graph.
• Defensive parseToolCalls (#111 stopgap). Three known
  malformations now repaired before the OpenAI shape returns: spec form,
  flat form, and XML-tag corruption. Token-attractor root cause
  (calibration retrain) deferred to a follow-up release.
• Daemon UX hardening. Gpu::init() failures convert from
  expect() panic to a friendly platform-specific checklist via
  report_gpu_init_failure() + exit(1). Bun stack on the CLI side
  also caught: Engine.recv() cleanly process.exit(code)s when the
  daemon early-exits, instead of throwing through a stack trace.
• gfx1152 / Strix Halo APU arch gating. Added to all RDNA 3.5
  dispatch lists. Does not yet fix #50 (segfault on --precompile);
  awaiting reporter backtrace.

Tooling

• scripts/speed-gate.sh DPM warmup (this release). bench_run
  now sets HIPFIRE_DPM_WARMUP_SECS=3 so pp32 measurements are
  reproducible regardless of GPU thermal state. Cold-DPM penalty was
  ~16% on the 32-token prefill probe; baseline 1240 was implicitly
  warm-captured and unreproducible across fresh-process runs without
  this fix.

Known caveats

• MQ3 collapses on sub-9B models. 0.8B and 4B in MQ3 are advisory
  only — they parse and dispatch but quality drops below MQ4 by a wide
  margin on real prompts. Matches QuIP# / sub-4-bit literature.
• MQ2 is refused by default. The quantizer requires
  --format mq2 --i-know-this-is-broken to opt in. Lloyd-Max MQ2
  (qt=19) and Lloyd-Max MQ3 (qt=20) are the path forward; spike
  shipped this cycle, full PRs to follow.
• MQ3 + MoE / A3B is unsupported. The MQ3 batched path lacks an
  MoE-branched WMMA kernel; daemon refuses MQ3 weights inside
  DeltaNetMoe / FullAttnMoe layers at load time.
• #111 token-attractor unresolved. The parser stopgap masks
  symptoms; calibration retrain is the real fix and lands in a
  follow-up release.
• #50 (gfx1152 segfault) still open pending reporter data.
• #119 (ROCm 7.2 / clang 22 regression on Strix Halo) filed this
  cycle; no engine-side fix yet.

Upgrade

hipfire update                      # if installed via curl-bash
# or
git pull && cargo install --path crates/engine

Windows: re-run install.ps1 — daemon.exe + kernel blobs refresh
automatically.

v0.1.8-alpha.2

Second cycle of contributor PRs landing in the same calendar day. Five PRs merged (#71, #72, #73, #74, #75), three Codex-flagged config-path hardening passes on the just-merged DDTree wire-up, plus a stale-path bugfix that was blocking 27B DFlash measurement in the speed-gate.

See CHANGELOG.md for full notes.

Highlights

• gfx12 WMMA dispatch is now feature-complete for HFQ4 prefill on RDNA4. PR #71 (@RobinVanCauter) closes the residual GEMM gap. R9700 numbers vs master: 9B prefill +29.7%/+31.3% (pp32/pp128), 27B prefill +27.8%/+42.4%.
• DDTree wire-up + Path C PRD (PR #72, @flamme-demon). Opt-in via HIPFIRE_DDTREE_BUDGET=<n>; default decode path bit-exact preserved. Local validation on 7900 XTX with Qwen3.6-27B + DFlash MQ4 draft: 12/12 attractor-clean across the path-c-smoke.sh --full battery, on the exact target/draft pair where Paths A/B1 single-token-attractor failed.
• gfx908 / MI100 CDNA1 bring-up (PR #67, @linus-amg). Wave64 dispatch on 4 fused projection sites + a new fused_gate_up_hfq4g256_wave64 kernel. Cross-process verified: 9B decode +9.3%, 4B decode +4.9%, A3B MoE decode +11.0% on real 2× MI100. New tests/speed-baselines/gfx908.txt.
• Opt-in HFQ4 MMQ prefill path (PR #73, @KotDath). Q8_1 + i8 WMMA over 128×128 tiles, gated to HIPFIRE_MMQ=1 on RDNA3/3.5. Targets the Strix Halo prefill gap vs llama.cpp (#60); +19.8% on 4B pp256 (gfx1100) once batch amortizes.

API

• /v1/chat/completions thinking-mode fix (#74). Per-model max_think_tokens was silently dropped on the OpenAI-compatible API path, causing empty message.content on thinking-mode models. Reproducer in #74. Also fixed prompt_tokens: 0 hardcode.

CLI

• hipfire list shows .mq6 models (PR #75, @Nereuxofficial).

Speculative decode hardening

• HFQ6 WMMA graph-capture safety — all 6 HFQ6 WMMA wrappers (3 gfx11 + 3 gfx12) migrated to graph-capture-safe launch path. Closed a hipGraph dangling-kernarg bug class on the 6-bit prefill path.
• DDTree daemon config hardening — three Codex-flagged crashable env-var paths in the just-merged DDTree wire-up: budget cap (256), topk cap (kernel-aligned at 8), Path_C value validation. Replaces silent OOM / silent fallback with clear stderr lines.

Tooling

• scripts/speed-gate.sh 27B DFlash draft path fix (#61). Reported by @m0n5t3r as MISSING_DRAFT despite the file being downloaded — gate was hardcoding the legacy filename. Now accepts both names.

Issues filed for follow-up

• #65 — gfx12 WMMA: tune 9B prefill (multi-row, K-tile, s_prefetch, launch_bounds). RDNA4 follow-up to PR #71. Hardware: R9700 / 9070 XT.
• #70 — gfx908 / MI100 CDNA1: port MFMA prefill kernels (4 kernels + channel-tests). Closes ~35× prefill gap vs gfx1100.
• #41 — DDTree on gfx1100 RoPE phase-skew. Superseded by PR #72's Path C orchestrator (different mechanism, attractor-clean). Closing in 7 days unless reopened.

Known issues

• #60 reporter (@h2252) hit a --gen 0 panic on bench_qwen35_mq4 pre-alpha.2. Cannot reproduce on master; many of today's commits could have addressed it incidentally. If you hit this on alpha.2, please retry with RUST_BACKTRACE=1 and post the trace on #60.
• #68 (Windows + qwen3.6:27b VL trace) — root-caused to the v0.1.0-alpha-pinned daemon.exe; alpha.1's fresh binary should resolve. Awaiting reporter confirmation.
• #50 (gfx1152 / Strix Halo APU segfault) — different SKU than the gfx1151 work that landed today. Awaiting bt + dmesg + cache-clean repro.

Upgrade

hipfire update                      # if installed via curl-bash
# or
git pull && cargo install --path crates/engine

Windows: re-run install.ps1. The dynamic-release-query (#69) will pull the fresh daemon.exe automatically; asset-id cache stamp prevents stale-binary preservation.

Windows binaries

All three (daemon.exe, infer.exe, run.exe) are attached as release assets, cross-compiled from Linux via x86_64-pc-windows-gnu against the v0.1.8-alpha.2 commit. hip-bridge does runtime LoadLibraryW of amdhip64.dll, so no link-time HIP dependency.

v0.1.8-alpha.1 — gfx1201 unblock + gfx1151 autodetect + GGUF import + docs rewrite

Point release rolling up the post-v0.1.8-alpha work. Two contributor PRs land in this cycle plus a feature on the input side and a docs nuke + rewrite.

Highlights

• RDNA4 / 9070 XT unblock end-to-end (#54). gfx1201 WMMA codegen crash resolved via dispatch fallback to dot2 (6e100c2); first canonical gfx12 WMMA scaffold (6924f2a) with C-output mapping derived from CK trait math; full validated 5-kernel + 6-channel-test contributor port (PR #56, @RobinVanCauter) hardware-tested on R9700 silicon. C-mapping acc[j] = C[8*(tid>>4) + j][tid & 15] validated. Public dispatch still routes gfx12 through dot2 fallback pending perf measurement (#57); the WMMA methods on Gpu are exposed for channel-tests now and ready to flip when numbers land.

• gfx1151 / Strix Halo autodetect fix (PR #59, @KotDath). KFD gfx_target_version 110501 was decoding to gfx11051 instead of gfx1151. Same refactor incidentally fixes a latent same-class bug for any arch with non-zero step bytes (100302 → gfx1003 was equally wrong before this PR). Hardware-validated on Ryzen AI Max+ 395 / Radeon 8060S. Speed-baseline contribution welcome at #61.

• GGUF → HFQ4 / MQ4 import (new feature). hipfire quantize <file.gguf> accepts any GGUF the engine can load (Q4_K_M / Q8_0 / Q4_0 / Q6_K / F16 / BF16 / F32 source quants) and re-quantizes to hipfire's native HFQ4-G256 (default for dense Llama / Mistral / older Qwen) or MQ4-G256 (FWHT-rotated, opt-in for Qwen 3.5+ family). Tensor names translated GGUF → safetensors style at write; tokenizer preserved verbatim under meta.gguf_meta so converted files are self-sufficient.

  hipfire quantize ./tinyllama.Q4_K_M.gguf --install --register tinyllama:1b-gguf
  hipfire run tinyllama:1b-gguf "..."

  Quality is lower than quantizing from full-precision safetensors (it's a double-quant roundtrip — raise to --format hf6 or --format mq6 if you have the disk space).

• dflash_mode default flipped to off (was auto). DFlash is now opt-in: hipfire config set dflash_mode auto re-enables the genre-conditional auto-routing. Per-genre measurements showed DFlash a clear win on code, modest on instruct, and a net loss on long-form prose. Default-on overpromised; default-off + opt-in matches the actual win surface.

• Docs nuke + rewrite. The 39-file docs/ tree (mix of canonical user docs and operational artifacts: agent prompts, daily standups, port plans, perf checkpoints) consolidated to 10 canonical pages — GETTING_STARTED / CLI / MODELS / QUANTIZE / CONFIG / SERVE / BENCHMARKS / ARCHITECTURE / QUANTIZATION / methodology/perf-benchmarking. README cut 371 → 89 lines. New top-level LICENSE file (was missing despite README and Cargo.toml declaring MIT).

• New hipfire-kernel-tuning agent skill. Codifies the empirical kernel-perf methodology from this repo's git log: 6-step workflow, levers catalog (multi-row, K-tile depth, wave64 port, s_prefetch_data, WMMA / MFMA, fused projections, ISA flags, rocBLAS fallback), cross-arch dispatch routing rules, and five worked case studies — wave64 CDNA3 port (+2× MI300X decode, 4105035), nontemporal-load fake-win revert (-13% caught only by clean-baseline bisect, 34eb024), gfx11 WMMA C-mapping silent corruption (~6 weeks before catch, b7ac66a), and others.

• 27B DFlash perf restored (9a2c667). PR #32 cleanup-dead-wmma-kernels removed kernels that were on the K4 / WMMA dispatch path for 27B verify-shape GEMMs. Restored via revert + cherry-pick. Empirical anchor: 27B-3.5 LRU code DFlash @ max=120 = 199 tok/s τ=10.36 (was: 95 tok/s in pre-revert state).

• Vision correctness (#23 / PR #35). G↔B channel transposition fixed in preprocessing; regression test pinned at crates/engine/tests/channel_order.rs.

Issues filed for follow-up

Hardware-gated or discussion-stage, cleanly punt to the next release cycle:

• #57 — gfx12 WMMA dispatch wiring + perf vs dot2 (R9700 / 9070 XT)
• #58 — multi-GPU support roadmap (PP first, TP follow-up)
• #60 — prefill scaling regression vs llama.cpp at pp≥512 on 9B+
• #61 — gfx1151 Strix Halo speed-baseline + perf bench

Upgrade

hipfire update

No config migration. ~/.hipfire/config.json from v0.1.8-alpha remains compatible. Re-enable DFlash explicitly:

hipfire config set dflash_mode auto

Full changelog: CHANGELOG.md

v0.1.8-alpha — DFlash prompt-shape +26.7% (inspired by Lucebox)

v0.1.8-alpha (2026-04-25)

Major perf cycle on the DFlash branch. Headline: Phase 1 prompt-shape
adaptation lifts 27B-3.5 DFlash by +26.7% on PEP-8-style code prompts
(median 161 → 199 tok/s, τ 8.07 → 10.36). Plain DFlash 10-prompt
HumanEval mean at n_gen=256 sits at 146.9 tok/s — alongside Lucebox's
ggml/CUDA published numbers on parallel hardware (RTX 3090). DFlash
work this cycle was substantially inspired by Lucebox
— credit to Davide Ciffa for published targets, n_gen-aware bench
methodology, and pointers at perf opportunities.

Features

• Prompt-shape adaptation (Phase 1): env-gated \n{3,} → \n\n
  collapse before tokenize. Eliminates rare BPE token 1358 (\n\n\n)
  in favor of HOT token 271 (\n\n) on Qwen3.5/3.6 vocab. Default
  OFF — opt-in via prompt_normalize=true (config TUI / CLI) or
  HIPFIRE_NORMALIZE_PROMPT=1. Impl: engine::tokenizer::maybe_normalize_prompt,
  10 unit tests, wired into 4 entry points (dflash_spec_demo, daemon,
  run, triattn_infer).
• Token heat diagnostic (Phase 2): HIPFIRE_PROMPT_TOKEN_HEAT=1
  triggers per-position BPE merge-rank heat dump at every encode site.
  HIPFIRE_PROMPT_HEAT_JSON=1 for machine-readable JSON to stdout.
  Standalone tool crates/engine/examples/encode_prompt.rs.
• EOT-stop fix: Tokenizer::is_terminator(id) + eot_id field.
  Daemon, run, dflash_spec_demo now stop on <|endoftext|> too, not
  just <|im_end|>. Fibonacci attractor loop in raw-text DFlash dead.
• DFlash drafts on HuggingFace: pullable via hipfire pull —
  qwen3.5:9b-draft, qwen3.5:27b-draft, qwen3.6:27b-draft. Plus
  qwen3.6:27b (15 GB target) added to registry. Files land at
  ~/.hipfire/models/<canonical-name> matching daemon auto-discovery.
• DDTree task #101 (tree-aware LA kernel): conv1d_silu_split_tree.hip
  • gated_delta_net_q8_tree.hip shipped with parent_indices plumbing
    through spec_step_ddtree_batched. Default ON post-validation.
• Adaptive-b continuous scheduler: with hysteresis + B_MAX clamp to
  draft-trained block size. Restores dflash_spec_demo's adaptive
  behavior in hipfire serve.

Tools / harnesses

• scripts/sweep_dflash_full.sh — unified 3 model × 2 mode × 3 genre
  bench harness (post-EOT-fix).
• scripts/bench_humaneval_dflash.py — HumanEval bench (Lucebox parity
  methodology, n_gen=256).
• benchmarks/prompts/lru_cache_pep8_strict.txt + lru_cache_single_blank.txt
  — md5-locked canonical bench prompts (CLAUDE.md mandate).
• crates/hipfire-quantize/src/bin/draft_to_mq4.rs — draft requantizer
  (+7.35% tok/s on 27B draft path).
• coherence-gate-dflash.sh is now the canonical correctness gate
  (replaces deprecated quality-gate.sh byte-exact baselines).

Kernel work

25 kernel files touched. Major:

• WMMA fast paths: gemm_hfq4g256_residual_wmma_{k2,k2x32,ksplit}.hip,
  gemm_qkv_hfq4g256_wmma.hip, gemm_qkvza_hfq4g256_wmma.hip,
  gemm_gate_up_hfq4g256_wmma.hip, gemm_mw16_residual_wmma.hip
  (tasks #73-#86): wo_residual 41% → grid-starvation fixed; gate_up
  50% → 75% BW; qkvza 46% → 70% BW; lm_head 32-rows/block.
• Tree-aware: conv1d_silu_split_tree.hip, gated_delta_net_q8_tree.hip.
• Fused: fused_rmsnorm_mq_rotate.hip, fused_qk_l2_norm_scale.hip,
  fused_sigmoid_alpha_gate.hip.
• Batched: embedding_hfq4g256_batched.hip, embedding_q8_batched.hip,
  kv_cache_write_q8_0_batched.hip, rope_partial_interleaved_batched.hip.
• 55+ hot-path hipMemset → hipMemsetAsync migrations (task #88).
• Async hipMemcpyHtoD for per-cycle uploads (task #89).
• wave_reduce intrinsic sweep (task #81).

CLI

• prompt_normalize: boolean config field — TUI toggle, per-model
  overridable, env-propagated through applyConfigEnv.
• Bonus: hipfire config set <bool-key> true|false parser fixed —
  was rejecting strings; now coerces "true"/"false" → bool. Incidentally
  fixes pre-existing dflash_adaptive_b, cask, experimental_budget_alert
  set paths.
• New aliases: qwen3.5-9b:draft, qwen3.5:9b:draft → qwen3.5:9b-draft.

Docs

• AGENTS.md (new) — tester-focused playbook complementing CLAUDE.md
  rules. Setup, smoke tests, gotchas table, flag reference.
• CLAUDE.md: "Prompt-structure τ sensitivity" hard rule (one newline
  = 17% τ swing) + "DFlash Coherence Gate" section.
• docs/methodology/perf-benchmarking.md: prompt-shape section + null-
  result log of attempts that looked like wins one-shell A/B but
  measured no-op on fresh probe.
• New PRDs: prompt-shape-adaptation.prd, task-93-path-c-trained-draft.prd,
  task-93-path-d-stale-context.prd, humaneval-2026-04-24-beats-3090.prd,
  prompt-structure-tau-discovery-2026-04-24.prd. Plus DDTree post-mortems.

Performance numbers (RX 7900 XTX, gfx1100, asym3 KV)

DFlash by genre (5-run medians, max=120, --no-chatml, normalize ON):

Model	code (HE/53)	code peak	prose (Rome)	instruct (sky)
27B-3.5	196.8 (4.45× AR)	218.6	49.6 (1.13×)	44.7 (~tie)
27B-3.6	185.5 (4.19× AR)	186.0	—	—
9B-3.5	329.1 (2.65× AR)	346.7	99.4 (0.79× ✗)	246.9 (1.99×)

DFlash is genre-conditional — net loss on 9B prose (-20%) because
draft-target argmax disagreement at high-entropy text. dflash_mode=auto
default handles this automatically for known-loss configs.

Lucebox parity (10-prompt HumanEval @ n_gen=256, plain DFlash linear):
146.9 tok/s mean — alongside Lucebox's published 112.82 (Chain DFlash)
and 135.80 (best DDTree b22 f16) on RTX 3090.

Roadmap (post-0.1.8)

GitHub issues:

• #38 Path D: Stale-context overlap pipelining (+5-15%)
• #39 Path C: Train custom DFlash draft (+15-30%)
• #40 Phase 3 prompt-shape: generalized rare-token rewriting (good first issue)
• #41 Fix DDTree on gfx1100: RoPE phase-delta skew
• #42 Mutable hipGraph: replace per-B verify cache
• #43 SSM intermediate persist-write
• #44 Vulkan backend MVP
• #45 v0.1.8+ Roadmap (living index) — pinned

Notes

• Master fixes merged into dflash for release: vision encoder F32 dequant
  fixes (#28), Qwen3.6-A3B HF upload (#34), batched-prefill determinism
  fixes (PR #28 follow-ups in master).
• Legacy byte-exact quality-gate.sh baselines are stale by design —
  branch-local artifacts, not the canonical gate. Coherence-gate-dflash
  is the gate going forward.
• DDTree on gfx1100 still has a structural RoPE phase-delta skew bug
  (commit 39aa358). Linear DFlash is the production path until issue
  #41 lands.

v0.1.7-alpha.3 — Qwen3.6-A3B fix + Phase 2 KV + coherence gate

Bundles 100 commits since v0.1.7-alpha.2. Points at merge commit `a60eda1` (PR #27).

Highlights

Qwen3.6-A3B no longer spirals — final-norm convention fix

Commit `1e01c0b`. The final `output_norm` on `qwen3_5_moe` (A3B, arch_id=6) is stored as a raw RMSNorm scale (mean ~+1.6), NOT as deviation-from-0 like the per-block norms. Our `load_norm_weight` unconditionally added 1.0, over-amplifying pre-lm_head activations by ~60% and tipping 3.6-MQ4 into infinite `` spirals on reasoning prompts. 3.5-A3B tolerated the over-amplification but was also subtly wrong.

Fix gated on `config.num_experts > 0`, so dense Qwen3.5 0.8B/4B/9B are untouched and byte-exact compatible with prior baselines.

Validated end-to-end: hermes-agent → SSH tunnel → MI300X serve → 3.6-A3B thinking:on closes `` and delivers answers instead of spiraling.

Qwen3.6:35b-a3b support (`cf3031f`)

Drop-in — same arch_id=6 as 3.5-A3B. Added REGISTRY entry + `qwen3.6` / `qwen3.6:a3b` aliases.

Phase 2 KV: physical_cap decoupled from max_seq

Commits `422fbf6`, `8f9c972`, `63ad050`. The advertised context window (`max_seq`) no longer forces VRAM allocation at `max_seq` × KV-bytes — the physical buffer is sized to `cask_budget + β + 256` and rolled via eviction. A3B can now advertise 4M+ tokens of context at ~22 GB VRAM on 24 GB cards.

CASK + DFlash eviction integration

Commits `1fbdcfa`, `adfe836`. DFlash speculative decoding now composes with TriAttention KV eviction. Downgrade guard: `cask=true` + `draft` logs a warning and falls back to plain TriAttention (m-fold + draft spec currently interacts pathologically).

rocBLAS MFMA path (task #130, CDNA3)

MFMA-accelerated prefill GEMMs on MI300X gfx942 (4.46× vs hand-rolled GEMV on A3B prefill). Kill-switch `HIPFIRE_ROCBLAS_OFF=1` for A/B benching or workaround. `HIPFIRE_ROCBLAS_ALL_ARCHS=1` opens the path to RDNA3 for smoke testing. Default-gated to CDNA3 arches.

Note: there's a separate rocBLAS + eviction prompt-KV corruption bug on MI300X (27B panics with `rocblas_gemm_ex hipError=6`, A3B produces `"The The The"` garbage under CASK). Worked around with `HIPFIRE_ROCBLAS_OFF=1` baked into the MI300X serve launcher. Real fix pending.

Coherence battery replaces byte-exact quality gate

New `scripts/coherence-gate.sh` runs a small (model × prompt) matrix through the daemon and writes a markdown report. Fails only on hard errors (daemon panic / zero tokens / timeout). Byte-exact `quality-gate.sh` (still available for manual use) was blocking legitimate numerical-correctness fixes (like this release's final-norm fix); the pre-commit hook now uses coherence. Committer (human or agent) reads the report and confirms coherence before the commit lands.

Smaller fixes

• `19f2c7b` — default A3B DFlash off (regression protection for reasoning tasks).
• `30022f0` — config TUI crash on missing meta for new 0.1.7 keys.

Upgrade

`hipfire update` pulls `origin/master` directly — no release-tag tracking — so this release is live for any user who runs it. Release tags are for GitHub discoverability.

Known issues (not fixed by this release)

• rocBLAS + eviction corruption on MI300X — workaround `HIPFIRE_ROCBLAS_OFF=1`.
• Local phrase repetition inside `` on 3.6-A3B (engine spiral is fixed; model still sometimes repeats a sentence 3–4× inside think block, which can demote the correct answer token via repeat_penalty → off-by-one final answer). Requires porting `apply_ngram_block` into the GPU `sample_top_p` kernel with scope reset at ``.

Cargo version note

`Cargo.toml` still reads `0.1.7-alpha` in this release's tree. Version string bump to match `0.1.7-alpha.3` deferred to the next commit. Does not affect runtime behavior or `hipfire update`.

🤖 Generated with Claude Code

v0.1.7-alpha: FlashTriAttn + CASK + A3B family + MI300X wave64

v0.1.7-alpha: FlashTriAttn + CASK + A3B family + MI300X wave64

Pre-release tag. Gated to full v0.1.7 on the outcome of the Hermes-agent
stack validation currently running on MI300X (task #125).

Highlights

• FlashTriAttn long-ctx wins shipped. DFlash speculative decode + TriAttention
  KV eviction composes cleanly. Measured on 7900 XTX / 9B MQ4 / ~1500-token
  prompt / 200-token decode / --cask-budget 512 --cask-beta 128:

  config	tok/s	τ
  DFlash baseline (full KV)	150	5.31
  FlashTriAttn (new 1M sidecar)	214 (+42%)	5.36 (−0.0)
  FlashCASK (m-folding, CPU merge)	124	2.83

  Earlier sidecars (small corpora) lost ~27% τ — that's gone now that cals are
  1M+ tokens. The FlashCASK τ drop is CPU-merge smoothing; GPU merge kernel
  (task #82) lands in 0.1.7 stable.

• Qwen3.5-35B-A3B and Qwen3.6-35B-A3B MoE end-to-end in DFlash. Batched
  MoE prefill + fused sigmoid+residual GEMV + indexed expert dispatch.

• MI300X (gfx942) wave64 port. 10 hot HFQ4 kernels re-written for
  block=[64,1,1] 2-rows-per-block. A3B decode 48.6 → 96 tok/s on MI300X
  (matches 7900 XTX baseline despite the 4× BW gap).

• Batched-prefill TriAttention tap — 4.5–5× faster sidecar calibrations.
  What made 1M-token cals across 5 targets feasible on one MI300X overnight.

• DFlash tape-replay rollback for multi-turn state recovery after verify
  mismatches without a full target re-run.

Bench snapshot (7900 XTX, MQ4)

DFlash τ + tok/s per prompt class (ctx=4K, no CASK):

model	short	code	math
4B	53 tok/s τ=1.27	92 tok/s τ=2.49	148 tok/s τ=6.0
9B	112 tok/s τ=1.52	461 tok/s τ=9.95	288 tok/s τ=5.77
27B	20 tok/s τ=2.21	41 tok/s τ=5.66	42 tok/s τ=6.14

5-model sidecar calibration (1M tokens wikitext):

model	mean r̄	notes
4B	0.564	dense
9B	0.629	dense, best generalization (tested 0.672 on code prompt)
27B	0.542	dense
3.5-A3B	0.552	MoE
3.6-A3B	0.552	MoE

Paper Figure 3 target is r̄ ≈ 0.5 — all five above it.

CLI + daemon config

Per-model config via hipfire config set <key> <value> or
~/.hipfire/per_model_config.json:

dflash_adaptive_b   boolean   default true
cask_sidecar        string    default ""        # path to .triattn.bin
cask                boolean   default false     # m-folding on top of sidecar
cask_budget         int       default 512
cask_beta           int       default 128
cask_core_frac      float     default 0.5
cask_fold_m         int       default 2

Daemon load message accepts these in params. cask_sidecar is accepted
and logged in alpha; the serve-time generate-loop integration lands in 0.1.7
stable. For now, use dflash_spec_demo directly for the --cask-sidecar path.

Pending for v0.1.7 stable

• Wire cask_sidecar + adaptive-B through the daemon generate loop.
• Hermes agent validation outcome (task #125) — gates stable release.
• GPU-side CASK merge kernel (task #82) for net-positive FlashCASK.
• DDTree integration into CLI/daemon (currently τ-positive, not yet tok/s-positive
  without hipGraph coverage).

Test plan

• Workspace cargo check clean on dflash branch
• cargo build --release --features deltanet --example daemon succeeds
• Local 7900 XTX bench sweep (scripts/dflash_branch_bench.sh) — all
  targets complete, τ + tok/s recorded
• 5 sidecars calibrated + rsynced back from MI300X
• triattn_validate --load-sidecar validates all 5 sidecars
• FlashTriAttn +42% speedup reproduced on 9B long-ctx
• Hermes agent + hipfire-daemon stack validation on MI300X (task #125)
• Quality gate sweep — NOTE: 0.1.6 baselines diverged from dflash output
  (known since branch start); user flagged as "fossil, ignore" 2026-04-17

🤖 Generated with Claude Code