hipfire

Fast local LLM inference for AMD GPUs.
Rust + HIP + Redline. No Python in the hot path. Ollama-style UX.

MQ4R + Redline · Models · GPU support · Install · Docs

hipfire pull qwen3.5:4b
hipfire serve qwen3.5:4b -d
hipfire chat qwen3.5:4b

One-shot inference uses the same model registry and serving stack:

hipfire run qwen3.5:4b "What is the capital of France?"

The daemon exposes an OpenAI-compatible API on 0.0.0.0:11435.

Current stable release: v0.2.1. The next release is v0.3.0, headlined by MQ4R and Redline across RDNA, and adding Qwen 3.8 27B and Muse Glimmer 30B. See CHANGELOG.md.

Curated weights are published through huggingface.co/hipfire-models and the per-model repositories recorded in the dynamic registry.

Discord: https://discord.gg/F3BaywB8Rs

MQ4R + Redline

MQ4R is the performance-oriented Qwen 3.6 35B-A3B SKU. It combines uniform MQ4 attention and gate-side weights with graded routed experts and the fused gate path.

hipfire pull qwen3.6:35b-a3b-mq4r
hipfire run qwen3.6:35b-a3b-mq4r \
  "Implement a bounded lock-free queue in Rust."

The model is 18.7 GB and requires approximately 22 GB of available VRAM. MQ4R prioritizes throughput; use the default MQ4P model or MFP4, MQ5, or MQ6 when quality matters more than maximum decode speed.

Redline is hipfire's in-tree dispatch and retained-replay substrate. It records the actual kernel graph, derives resource dependencies, retains invariant command state, and lowers validated paths through public ROCr queue interfaces.

Redline supports the RDNA architecture family from RDNA1 through RDNA4. Optimized routes remain architecture- and workload-specific. Unsupported graphs, failed shadow validation, ABI mismatches, queue faults, and model changes fail closed to ordinary HIP dispatch.

Qwen 3.6 35B-A3B MQ4R performance

Ordinary autoregressive decode with Q8 KV. No MTP, DFlash, speculative acceptance, reduced-output benchmark, or manual clock pinning.

Short-context measured ranges:

The eight-turn column is the mean decode speed across the user-facing multi-turn serving run. Final-turn context is reported explicitly because the three recorded sessions did not all terminate at exactly 22K.

The gfx1201 campaign raised ordinary MQ4R autoregressive decode from approximately 110 tok/s to 203.9 tok/s. See the gfx1201 campaign report and Redline integration boundary. To reproduce the sealed TG128 route, model, and sampling fixture on a supported GPU, run python3 -m tools.redline golden; the Golden Redline guide also covers making the validated fixture the serve default and connecting Hermes, Pi, or another OpenAI-compatible client.

Curated model registry

The registry currently contains 61 pullable model entries. Run hipfire list -r to see the authoritative live list.

Common aliases include qwen3.5, qwen3.6, qwen3.8, qwen3, carnice, qwopus, deepseek4, deepseek-v4, muse-glimmer, and vibethinker.

Carnice uses the Hermes tool-call format. Plain Qwen 3.5 and 3.6 use their native Qwen XML tool-call format.

See docs/MODELS.md for sizes, minimum VRAM, recommended sampling settings, sidecars, artifact provenance, and bring-your-own-model flows through hipfire quantize.

GPU support

Architecture-specific kernels are selected through typed dispatch tables. Unsupported specializations return to the correct portable or architecture fallback instead of being applied to neighboring GPU families.

Why

AMD GPUs are capable inference devices, but tuning and runtime support vary widely across consumer, professional, APU, and datacenter products.

hipfire supplies its own Rust runtime, model implementations, quantization formats, dispatch layer, and HIP kernels. ROCm is loaded dynamically; there is no Python, PyTorch, CUDA translation layer, or third-party inference engine in the hot path.

Redline removes launch overhead after the ordinary graph has been proven safe. It does not replace the model implementation or bypass hipfire's correctness gates.

Performance snapshots

Historical — 7900 XTX (gfx1100)

Historical decode snapshot measured with the then-default asym3 KV configuration (FlashAttention auto):

Historical DFlash measurements using the legacy asym3 / max=120 method reached 218 tok/s peak on 27B HumanEval/53 (4.45× over AR) and 372 tok/s peak on 9B. These numbers are retained for historical context; they are not current performance baselines. Current DFlash claims use q8 / max=256 with prompt and binary hashes recorded. See docs/BENCHMARKS.md for the full per-genre table and methodology notice.

RDNA4 (gfx1201, Radeon AI PRO R9700)

The DeepSeek V4 Flash rows are n=3 fresh-process medians on deepseek-v4-flash-0731.mq2r (sha256 cbf2bbcf…), greedy, speculation off, --kv f32, 2052-token prompt. MQ2R is the current deepseek-v4-flash SKU and roughly doubles the superseded MQ2-Lloyd row below it; benchmark against MQ2R, not the Lloyd figure. TP3 trades ~2% decode for ~24% faster prefill (481 vs 389 tok/s) and leaves a fourth card free.

Experimental long-context compression and eviction are opt-in. PFlash is off by default, TriAttention sidecars do not auto-attach, and CASK m-folding is disabled. Generate a sidecar with hipfire sidecar-gen <model>, then set memory.cask.sidecar to its exact path (or explicitly enable memory.cask.auto_attach). Set memory.cask.enabled=true only when m-folding is intended. See CONFIG.md for details.

Install

Linux with ROCm 6 or newer:

curl -fsSL https://raw.githubusercontent.com/warpfront/hipfire/master/scripts/install.sh | bash

# Or install the integration branch for testing:
curl -fsSL https://raw.githubusercontent.com/warpfront/hipfire/master/scripts/install.sh \
  | bash -s -- --branch beta

RDNA4 requires ROCm 6.4 or newer. gfx1151 requires ROCm 7.2 or newer. Run hipfire --version for a concise build ID or hipfire version to compare the installed binary, managed source checkout, and daemon. Managed Linux installs can switch revisions with hipfire update @beta, hipfire update --tag v0.2.1, or hipfire update --commit <sha>.

To uninstall a managed Linux install while keeping downloaded models and settings:

curl -fsSL https://raw.githubusercontent.com/warpfront/hipfire/master/scripts/uninstall.sh | bash

Add --dry-run to preview it. --purge also deletes all data under ~/.hipfire and requires explicit confirmation.

For Windows, source builds, and verifying the install: docs/GETTING_STARTED.md.

NixOS

First-class support via Nix flake. See docs/NIXOS.md.

nix develop github:warpfront/hipfire  # dev shell with Rust + ROCm
nix build github:warpfront/hipfire    # build package

NixOS module:

{
  inputs.hipfire.url = "github:warpfront/hipfire";
  # then in configuration.nix:
  services.hipfire.enable = true;
  services.hipfire.gpuTargets = [ "gfx1100" ];
}

Container (podman/docker)

A multi-stage Containerfile builds a slim deliverable inference image and a full-toolchain GPU gate-runner for reproducible PR/dev-build validation. See docs/CONTAINER.md.

podman build -f Containerfile --target runtime -t hipfire .
podman run --rm -it --device /dev/kfd --device /dev/dri \
  --group-add keep-groups --security-opt seccomp=unconfined \
  -v hipfire-models:/root/.hipfire/models \
  -v hipfire-kcache:/var/cache/hipfire \
  hipfire run qwen3.5:4b "2+2="

Inspiration: Lucebox

hipfire's DFlash work was substantially shaped by Davide Ciffa's Lucebox DFlash on ggml — a standalone C++/ggml/CUDA DFlash for Qwen 3.5-27B on a single RTX 3090. Different stack, different vendor — but Lucebox's blog gave us concrete published numbers to target, n_gen-aware bench methodology, and pointers at where the fat is. Cached snapshot at .research-cache/lucebox-dflash27b.html for forensic reproducibility.

Inspiration: gfx906 (MI50/MI60) optimizations

hipfire's gfx906 prefill MMQ kernel and AR-decode optimizations were shaped by two community forks of llama.cpp that target Vega 20:

• iacopPBK/llama.cpp-gfx906 — the original fork that ported and tuned gfx906-specific code paths (warp-cooperative GEMV via half-wave split, Y-tile prefetch via inline-asm global_load_dword, __builtin_amdgcn_readfirstlane-based SGPR hoisting, separate HBM-load → register-cache → LDS-store pipelining in the MMQ body). The "2602.01 version" commit eec153c086df6a9e7a69499bea3639597c085fff was the canonical reference we audited against.
• skyne98/llama.cpp-gfx906 — fork-of-fork that propagates iacop's optimizations (commit 42c298c "port iacop optimizations") and tracks upstream more aggressively. The accompanying skyne98/wiki-gfx906 is the best public reference for gfx906 ISA quirks (LDS bank-conflict patterns at stride 32, dp4a issue-rate ceiling, Q8_1 activation layout) — we used it as a sanity-check for several PMC-driven redesign decisions.

And of course an extra shout-out to ggml-org/llama.cpp itself: the templated mmq_x body in mul_mat_q.cu was the architectural scaffold we ported to gfx906 (templated mmq_x ladder, per-thread accumulator layout, MMQ_TILE_NE_K=32 sub-block factoring, Q8_1 quantize math). The inner loop is gfx906-specific; the outer shape is descendant.

A standalone gfx906 perf investigation log is at docs/perf-checkpoints/2026-05-05-gfx906-decode-investigation.md; the prefill MMQ redesign log is at docs/perf-checkpoints/2026-05-05-gfx906-mmq-redesign-final.md.

Documentation

License

hipfire is licensed under Apache-2.0 as of v0.3.0. See LICENSE, LICENSE-APACHE, and NOTICE for details. Releases up to and including v0.2.1 were dual MIT/Apache-2.0; that grant is not revoked.

Individual files whose substantive authors have not elected Apache-2.0 remain MIT-licensed (see LICENSE-MIT) and are identified by their per-file SPDX-License-Identifier. No contributor's file was relicensed in absentia. New contributions default to Apache-2.0 via DCO sign-off. See CONTRIBUTING.md for the contributor side and docs/governance/relicense-2026-05.md for the decision record (including the 2026-05-19 course correction and the v0.3.0 move to outbound Apache-2.0).

Original architectural innovations originating in hipfire are catalogued in PRIOR-ART.md; derivative works (including reimplementations informed by hipfire's design) should attribute the corresponding inventions per AGENTS.md.

Contributing

See CONTRIBUTING.md. Install local hooks with ./scripts/install-hooks.sh. The no-GPU CI subset is ./scripts/no-gpu-ci.sh; it does not replace the hardware gates. There is no single canonical correctness gate — the retired scripts/coherence-gate*.sh batteries are not acceptance evidence and coherence-gate-dflash.sh no longer exists. Select the route your change owes from docs/VALIDATION.md: scripts/serve_harness.py for generation and state-lifecycle changes, scripts/redline_daemon_harness.py for kernel, dispatch, graph, or Redline-replay changes. Architecture-specific work must also pass its channel test, anti-bleed checks, and relevant speed gate. Don't bypass the gates with --no-verify — see methodology/perf-benchmarking.md.