{
"markdown": "\n\n# fucina\n\n### Gemma 4 and Qwen3 / Qwen3.5 / Qwen3.6, forged for the NVIDIA DGX Spark\n\n***fucina*** *— Italian for **forge**: the smithy where raw model weights are hammered into a fast engine for one machine.*\n\nA from-scratch inference engine, hand-tuned for exactly one accelerator — the **DGX Spark GB10**\n(Blackwell, `sm_121a`, CUDA 13). It serves **Gemma 4 12B**, **Qwen3 / Qwen3.5 / Qwen3.6** (dense\nand MoE, GGUF / FP8-safetensors / NVFP4), and the experimental **DiffusionGemma 26B-A4B** text-\ndiffusion MoE — FP8/NVFP4 Tensor-Core attention, CUDA-graph decode, continuous batching over a\npaged KV cache, speculative decoding, and OpenAI/Anthropic-compatible APIs, all in a single static\nbinary.\n\n[Features](#-features) · [Quick start](#-quick-start) · [Model support](#-models) ·\n[Performance](#-performance) · [Continuous batching](#-continuous-batching--paged-kv) ·\n[Speculative decoding](#-speculative-decoding) · [HTTP API](#-http-api) ·\n[DiffusionGemma](#-diffusiongemma) · [OKF knowledge bundle](docs/okf/)\n\n\n-76B900)\n\n\n\n\n\n\n
\n\n> [!WARNING]\n> **fucina is experimental, hardware-specific, and provided with no support.** It is built and\n> tested for a single accelerator — the NVIDIA DGX Spark GB10 — with a toolchain expected at\n> `/usr/local/cuda-13` and `/usr/local/go`. It is **not portable** to other GPUs as-is. Licensed\n> **Apache-2.0** and shipped **as-is, with no warranty and no support**. See\n> [Hardware support](#-hardware-support) and [Status](#-status).\n\n---\n\n## About\n\n**fucina** runs Google's **Gemma 4 12B** and Alibaba's **Qwen3 / Qwen3.5 / Qwen3.6** families\n(dense and sparse-MoE) entirely on the GPU and serves them over OpenAI- and Anthropic-compatible\nHTTP APIs, plus one-shot and interactive CLI modes. It started as a focused experiment in *how\nfast a single Blackwell GB10 can drive a dense 12B model* and has since grown into a small multi-architecture\nengine — but it still bets everything on one accelerator instead of portability: FP8/NVFP4 Tensor\nCores, position-independent CUDA-graph decode, an on-GPU sampler, continuous batching over a paged\nKV cache, and speculative decoding measured head-to-head against `llama.cpp` and `vLLM`.\n\nThere is no model-selection flag: `-m ` is enough — the architecture (Gemma-4 / Qwen3 dense /\nQwen3 MoE / Qwen3.5-3.6 hybrid dense / Qwen3.5-3.6 MoE) and the weight format (GGUF quant, official\nFP8-block safetensors, or NVFP4 — generic or NVIDIA ModelOpt) are detected from the file itself. See\n[Models](#-models) for the full support matrix and exact download commands.\n\nThe same binary also runs **DiffusionGemma 26B-A4B** — a block text-diffusion MoE model — through a\nseparate CUDA engine selected with `-dm`. See [DiffusionGemma](#-diffusiongemma).\n\n**Gemma 4 12B:** 48 layers (40 sliding-window + 8 global attention), GeGLU FFN\n(`gelu_pytorch_tanh`, intermediate 15360), RoPE / p-RoPE, logit softcap 30.0, vocab 262144, hidden\n3840, and a Q6_K tied LM head — loaded from **Q4_0 (QAT)** or **Q8_0** GGUF weights, or native\nNVFP4 safetensors.\n\n**Qwen3.5 / Qwen3.6 hybrid:** a Gated-DeltaNet (GDN) linear-attention mixer on most layers with a\nperiodic full softmax-GQA layer (`qwen35.full_attention_interval`), served dense (9B/27B) or as a\nsparse MoE (35B-A3B, top-k routed experts), from GGUF or official FP8-block / NVIDIA ModelOpt\nsafetensors. \"Qwen3.6\" checkpoints reuse this same hybrid architecture and detector — there is no\nseparate Qwen3.6 code path.\n\n---\n\n## ✨ Features\n\n**fucina is fast with:**\n\n- ⚡ **CUDA-graph decode** — single-token decode *and* the K-row batched speculative-verify forward\n are captured as **position-independent** graphs (device-resident position, KV writes *inside* the\n graph) and replayed each step, eliminating per-kernel launch overhead. Pre-captured at startup.\n- 🎯 **MTP speculative decoding** — prompt-lookup speculation (free) plus an optional **MTP draft\n head** (`--assistant`) that drafts novel text; one batched weight pass verifies many tokens at the\n exact target distribution. **>2× dense decode** at typical acceptance.\n- 🧮 **FP8 Tensor-Core attention** + **FP8 E4M3 KV cache** (1 byte/element) — half the KV bandwidth\n and a flat decode curve as context grows. The sliding-window layers use a **capped ring buffer**,\n so KV memory stays nearly context-independent (~1.5 GiB sliding regardless of `--ctx`). An optional\n **NVFP4 KV codec** (`FUCINA_KV_NVFP4=1`, native Blackwell FP4) trades a small accuracy cost for a\n 1.78× smaller KV footprint — memory-only, opt-in (see [`docs/kv-quant-exploration.md`](docs/kv-quant-exploration.md)).\n- 🧵 **Continuous batching over a paged KV cache** — independent sequences decode in one batched\n pass (vLLM-style block tables + free-list, split-K paged attention bit-identical to the contiguous\n path, CUDA graph per batch size). **Opt-in for Gemma-4** (`--batch`, implies `--paged-kv`) — the\n default single-flight path is unchanged there; **mandatory and auto-enabled for every Qwen3\n checkpoint** (Qwen has no single-flight path). See [`docs/continuous-batching.md`](docs/continuous-batching.md).\n- 🧑🤝🧑 **Multi-architecture, no model-select flag** — Gemma-4, Qwen3 dense, Qwen3 MoE, and the\n Qwen3.5/3.6 Gated-DeltaNet hybrid (dense and MoE) are all detected from the GGUF metadata or\n `config.json`, and served through the same binary and (mostly) the same code paths.\n- 📦 **Native quantized GEMV/GEMM** — Q4_0 / Q6_K / Q8_0 / Q4_K read directly via `dp4a`, with an\n optional repacked-Q4_0 coalesced-load decode path (Gemma) and a default Q4_K requant of the\n attention/GDN mixer + FFN/expert weights for Qwen3.5/3.6 dense **and** MoE checkpoints (smaller\n resident weight set, faster decode; `FUCINA_MOE_FP8=1` reverts to pure FP8). No BF16 materialize\n on the decode hot path.\n- 🧠 **On-GPU sampling** — the next token is selected on the device; no full-vocab logit copy back\n to the host when no repeat penalty is set.\n\n**fucina gives you:**\n\n- 🔁 **Prefix-reuse KV cache** — instead of re-prefilling the whole prompt each request, the server\n rewinds the single physical KV cache to the longest common prefix and prefills only the divergent\n suffix — the difference between sub-second and multi-second agentic turns. Rewinds stay exact within\n the sliding ring's window (covers same-conversation turns and speculation); a deeper divergence\n falls back to a full re-prefill (see `FUCINA_SLIDING_RING`).\n- 🌐 **OpenAI- and Anthropic-compatible APIs** — `/v1/chat/completions` and `/v1/messages`\n (streaming + non-streaming), `/v1/models`, `/health`, `/metrics`.\n- 🛠️ **Tool calling**, OpenAI-shaped on the wire for both dialects — Gemma-4's native format and\n Qwen3-Coder's XML `` form (auto-selected with the chat dialect, no flag) — plus a\n **thinking/reasoning channel** on both (reasoning as `reasoning_content`, controllable per request\n via `reasoning_effort` or `thinking`).\n- 💻 **Three run modes** from one binary — server, one-shot prompt, interactive REPL.\n- 📊 **Live observability** — `/metrics` exposes prefix-cache hit rate, prefill/decode throughput,\n and a `speculation` block (`accept_rate`, `tokens_per_forward`).\n\n---\n\n## 🚀 Quick start\n\n### 1. Install dependencies\n\nfucina runs on a **DGX Spark GB10** (see [Hardware support](#-hardware-support)). You need:\n\n| Dependency | Version | How |\n|---|---|---|\n| **CUDA Toolkit** | **13.0** | Expected at `/usr/local/cuda-13`; else pass `NVCC=`/`CUDA_HOME=` to `make`. [NVIDIA CUDA downloads](https://developer.nvidia.com/cuda-downloads) |\n| **Go** | **1.26** | Expected at `/usr/local/go`; else pass `GO=` to `make`. [go.dev/dl](https://go.dev/dl/) |\n| **Build tools** | gcc/g++, make, binutils | `sudo apt-get install build-essential` (DGX OS / Ubuntu) |\n| **huggingface_hub** | Python 3.10+ | To download GGUF weights: `pip install -U \"huggingface_hub[cli]\"` |\n| **CUTLASS** *(DiffusionGemma only)* | sm_120-capable headers | `pip install flashinfer` vendors them, or `git clone https://github.com/NVIDIA/cutlass`; pass `CUTLASS_DIR=` to `make` |\n\n> [!NOTE]\n> The dense Gemma 4 engine needs only CUDA + Go + build tools. **CUTLASS is required only for the\n> optional [DiffusionGemma](#-diffusiongemma) engine** — skip it if you only want the 12B.\n\n### 2. Build\n\n```sh\nmake\n```\n\n`make` compiles the CUDA static libraries (`nvcc -arch=sm_121a` → `cuda/libfucina.a` for the dense\nengine and `cuda/libdg.a` for the [DiffusionGemma](#-diffusiongemma) engine) and then the Go binary,\nverifying the cubin arch and the device-upload code path along the way.\n\n> [!NOTE]\n> Requires the [supported toolchain](#requirements): CUDA 13.0 at `/usr/local/cuda-13` and Go 1.26\n> at `/usr/local/go`, on a DGX Spark GB10.\n\n#### Build configuration (no hardcoded paths)\n\nAll machine-specific locations are overridable — nothing personal is baked in. Pass them as `make`\nvariables (or environment variables for the helper scripts):\n\n| Setting | Used by | Default | Purpose |\n|---------|---------|---------|---------|\n| `NVCC`, `CUDA_HOME` | `make` | `/usr/local/cuda-13` | CUDA 13 toolchain location |\n| `GO` | `make` | `/usr/local/go/bin/go` | Go 1.26 toolchain |\n| `CUTLASS_DIR` | `make` (diffusion) | `/path/to/cutlass` | CUTLASS include dir for the NVFP4 MoE GEMM |\n| `DG_GGUF` | `make` (diffusion targets) · `scripts/dg_dump_tensor.py` (env) | `./models/diffusiongemma-26B-A4B-it-Q4_K_M.gguf` | DiffusionGemma GGUF |\n| `DG_NVFP4_CKPT` | `scripts/dg_nvfp4_convert.py` (env / `--ckpt`) | — (required) | NVFP4 safetensors snapshot dir |\n| `LLAMA_GGUF_PY` | `scripts/dg_dump_tensor.py` (env) | auto (only if `gguf` isn't importable) | path to llama.cpp's `gguf-py` |\n\n```sh\n# Example: build with your own CUTLASS and model locations\nmake CUTLASS_DIR=/opt/cutlass DG_GGUF=/data/diffusiongemma.gguf\n\n# Helper scripts read env vars / flags — no editing required\nDG_NVFP4_CKPT=/data/dg-nvfp4-snapshot python3 scripts/dg_nvfp4_convert.py --inspect\nDG_GGUF=/data/dg.gguf python3 scripts/dg_dump_tensor.py /tmp/out\n```\n\n### 3. Download a model\n\n```sh\n# Gemma 4 12B — Q4_0 QAT (recommended; official Google QAT GGUF)\nhf download google/gemma-4-12B-it-qat-q4_0-gguf \\\n gemma-4-12b-it-qat-q4_0.gguf --local-dir ./models\n\n# Optional: the MTP draft head for faster decode (see step 4 / MTP)\nhf download unsloth/gemma-4-12b-it-GGUF \\\n MTP/gemma-4-12b-it-Q8_0-MTP.gguf --local-dir ./models\n\n# — or — Qwen3.6-35B-A3B, official FP8 (MoE; --local-dir gives a self-contained checkpoint dir,\n# no --tokenizer needed — see \"HuggingFace checkpoint paths\" under Models)\nhf download Qwen/Qwen3.6-35B-A3B-FP8 --local-dir ./models/Qwen3.6-35B-A3B-FP8\n```\n\nSee [Models](#-models) for the Q8_0 dense build, every Qwen3/3.5/3.6 checkpoint, and the\nDiffusionGemma weights.\n\n### 4. Run\n\n**As an OpenAI-compatible server:**\n\n```sh\nfucina -m ./models/gemma-4-12b-it-qat-q4_0.gguf --ctx 32768 --host 0.0.0.0 --port 8080\n# — or — a Qwen checkpoint (continuous batching auto-enables itself; no extra flags needed):\nfucina -m ./models/Qwen3.6-35B-A3B-FP8 --host 0.0.0.0 --port 8080\n```\n\n```sh\n# then, from another shell:\ncurl http://localhost:8080/v1/chat/completions \\\n -d '{\"messages\":[{\"role\":\"user\",\"content\":\"Hello\"}],\"stream\":true}'\n```\n\nSee [HTTP API](#-http-api) for every endpoint and a tool-calling example.\n\n**As an interactive REPL:**\n\n```sh\nfucina -m ./models/gemma-4-12b-it-qat-q4_0.gguf --interactive\n```\n\nA multi-turn chat with prefix-reuse caching. Commands: `/thinking LEVEL` (set the reasoning\nchannel — `off`/`on`/`low`/`medium`/`high`/`xhigh`; the thought channel renders dimmed and is\nbudget-bounded), `/reset` (clear the conversation), `/stats` (KV-cache hit rate), `/quit` (or\nCtrl-D). The REPL applies the chat template for the detected dialect (Gemma-4 or Qwen ChatML — no\nflag) and honours `--thinking` / `--repeat-penalty`, matching the server. Qwen checkpoints route\nthrough the same paged multi-sequence path as the server (single-flight is Gemma-only).\n\n**Qwen3.5 J-space debugging (experimental):**\n\nFucina can record a fitted Anthropic Jacobian lens for each generated token and can steer the\nresidual stream interactively. This path is intentionally slow and must never be enabled in a\nproduction server. The lens must match the model family/checkpoint and have the same hidden width/layer count.\nFor Qwen3.5-9B Base, a converged 458-prompt lens is published in\n`neuronpedia/jacobian-lens`; the replication guide below gives the exact path and conversion.\n\n```sh\n# FP8 inference operators have no autograd formula. Materialize an offline BF16 fitting copy;\n# fucina itself should continue using the original FP8 checkpoint.\npython scripts/dequantize_fp8_hf.py /path/to/Qwen3.5-FP8 /tmp/Qwen3.5-BF16-for-jlens\n\n# Fit selected workspace layers and emit both .pt and .fjls. Run this with PyTorch,\n# transformers, accelerate, and anthropic/jacobian-lens installed; fitting is expensive.\npython scripts/fit_fucina_jlens.py \\\n --model /tmp/Qwen3.5-BF16-for-jlens \\\n --output qwen35_jacobian_lens.pt \\\n --layers 8,12,16,20,24,28 --n-prompts 8\n\n# Or convert a lens that was already fitted with the upstream package.\npython scripts/convert_jlens.py qwen35_jacobian_lens.pt qwen35.fjls --model-layers 32\n\n# Trace + permit interventions. --jspace alone traces but disables steering commands.\nfucina -m ./models/Qwen3.5.gguf --jspace-debug --j-lens qwen35.fjls \\\n --jspace-out /tmp/qwen35-jspace.jsonl --jspace-top-k 8\n```\n\nThe JSONL contains one record per generated token with per-layer decoded strings and\nfull-vocabulary probabilities. It records both `source_token` (the residual position being read)\nand `sampled_token` (the next token predicted from it), avoiding the common causal off-by-one\nmistake in autoregressive traces. Numeric ids are retained as replay metadata. Interactive commands:\n\n```text\n/jdump print the latest fitted-layer readouts\n/jsteer \" Paris\" 0.15 all steer toward a token word at all fitted layers\n/jsteer \" violence\" -0.10 8,16,24 suppress a token word at selected layers\n/jclear disable steering\n```\n\nSteering follows the J-Lens direction `J_lᵀ·w_token`, normalized and scaled by the current\nresidual norm. Strength is clamped to `[-1,1]`. Quote token words to preserve their significant\nleading whitespace. Numeric ids remain accepted only as an unambiguous fallback when multiple\nvocabulary entries decode to the same visible text.\n\nA runnable replication of the paper's Figure 9 citrus directed-modulation example is included:\n\n```sh\npython scripts/workspace_citrus_example.py \\\n --model ./models/Qwen3.5.gguf \\\n --lens ./qwen35.fjls \\\n --top-k 20 \\\n --output-dir workspace-citrus-run\n```\n\nIt runs a no-side-task control and a `concentrate on citrus fruits` condition while the model\ncopies `The old painting hung crookedly on the wall.` It then compares decoded J-space words at\nthe generated token covering `ook` in `crookedly`, writing both traces and `summary.json`. This is\na qualitative replication on a different model, not a reproduction of Anthropic's Sonnet 4.5\nnumbers. See [`docs/workspace-paper-replication.md`](docs/workspace-paper-replication.md) for the\nmethod, causal trace convention, and scientific caveats.\n\n**As a one-shot prompt:**\n\n```sh\nfucina -m ./models/gemma-4-12b-it-qat-q4_0.gguf -p \"Write a haiku about CUDA.\" -n 100\n```\n\n**With MTP speculative decoding** (Gemma-4 only) — add the draft head for faster decode; works in\n*all three* modes (here, the REPL):\n\n```sh\nfucina -m ./models/gemma-4-12b-it-qat-q4_0.gguf \\\n --assistant ./models/MTP/gemma-4-12b-it-Q8_0-MTP.gguf --interactive\n```\n\nQwen checkpoints get free-standing **prompt-lookup** speculation automatically on dense models\n(not MoE) — no `--assistant` head applies there. See [Speculative decoding](#-speculative-decoding)\nfor how it works and how to observe acceptance live.\n\n---\n\n## 📈 Performance\n\n**Gemma-4 12B**, measured against `llama.cpp` (`llama-server`) on a fair side-by-side harness\n([`scripts/pi_bench.py`](scripts/pi_bench.py): identical transcript, temperature 0, thinking on,\nMTP draft head on both engines):\n\n| Phase | Result vs llama.cpp |\n|-------|---------------------|\n| **Decode** | **Parity-to-ahead overall; +15–20% at high context (≥5k tokens)** — throughput *rises* with context as acceptance climbs and FP8-KV attention scales, the regime that dominates long agentic sessions. |\n| **Prefill** | **Steady-state tied**; the residual gap is one-time cold turns (small-N GEMM efficiency), not steady throughput. |\n| **Tool calling** | Matches on the easy suite, ahead on hard agentic scenarios, at equal-or-faster latency. |\n\nThe decode wins come from **CUDA-graph launch-bubble removal and speculation acceptance (τ)** — not\nweight-load width. Single-token decode on GB10 is **bandwidth-bound on total weight bytes**, so\nwider (128-bit) loads change instruction count but not bytes and do not help. Watch the\n`speculation` block in `/metrics` (`accept_rate`, `tokens_per_forward`) to observe τ live.\n\n### Measured throughput & memory\n\nSingle-stream decode on the **GB10** (`sm_121a`, CUDA 13, temperature 0). Decode is the\nsteady-state, bandwidth-bound metric; resident is the weight footprint (KV cache adds the\ncontext-dependent amount in the note above).\n\n| Model | Format | Decode (tok/s) | Weights resident |\n|-------|--------|---------------:|-----------------:|\n| **Gemma-4-12B-it** | Q4_0 (QAT) | ~28 base · **~57** with MTP | ~8.0 GB¹ |\n| **Gemma-4-E4B-it** | Q4_0 (QAT), native dp4a MMVQ | **45.2** | **7.05 GB** |\n\n¹ 6.96 GB packed-Q4_0 decode copy + 1.0 GB (Q6_K→Q8_0 head); `FUCINA_NO_PACKED=1` drops ~7 GB.\n\n**Gemma-4-E4B decode path** — same checkpoint, decode speed is purely a function of the weight\nbytes read per token. The native Q4_0 path reads the original QAT nibbles directly via the shared\ndp4a MMVQ kernels (the same ones the 12B uses), **2.6× faster than BF16 decode at ~half the memory**:\n\n| E4B decode path | Decode (tok/s) | Weights resident @ 4096 ctx |\n|-----------------|---------------:|----------------------------:|\n| BF16 (dequantized at load) | 17.6 | ~12.5 GB |\n| NVFP4 hybrid (re-quant) | 38.6 | 15.4 GB |\n| **native Q4_0 (dp4a MMVQ)** | **45.2** | **7.05 GB** |\n\n**Qwen3.5 vs vLLM — official position (2026-07-18, fresh contemporaneous head-to-head,\n`bench_serving.py` per [`benchmark-evidence/PROTOCOL.md`](benchmark-evidence/PROTOCOL.md); full record in\n[`docs/sota-gb10-qwen3-mission-plan.md`](docs/sota-gb10-qwen3-mission-plan.md)):**\n\n**fucina wins 11 of 12 aggregate-throughput concurrency cells while providing\nbyte-identical run-to-run determinism that vLLM does not offer.**\n\n**Qwen3.5-35B-A3B-FP8 (MoE) — all 6 cells won** (agg tok/s):\n\n| N | 1 | 2 | 4 | 8 | 16 | 32 |\n|---|---|---|---|---|---|---|\n| fucina | 59.0 | 101.4 | 134.0 | 229.8 | 320.1 | **472.4** |\n| vLLM | 14.0* | 74.1 | 111.7 | 155.2 | 207.2 | 321.3 |\n\nThe 11/12 claim is aggregate throughput, not every burst-TTFT cell. In the same fresh\n2026-07-18 sweep, MoE N=32 TTFT was 670/722 ms (med/p95) vs vLLM 312/316; fucina's\nsingle 3,500-token TTFT remained ahead (4,367 vs 5,837 ms).\n\n**Qwen3.5-9B-FP8 (dense) — 5 of 6 cells won** (agg tok/s):\n\n| N | 2 | 4 | 8 | 16 | 32 |\n|---|---|---|---|---|---|\n| fucina | 59.3 | 117.3 | 204.6 | 313.1 | 438.8 |\n| vLLM | 44.2 | 85.8 | 164.4 | 280.8 | **521.8** |\n\nThe dense N=32 cell (−16%) is a **measured, triply-proven kernel-class boundary**\n(tensor-core MMA vs dp4a GEMV — see [`docs/qwen35-d32b.md`](docs/qwen35-d32b.md)): closing it\nwould require a reduction order that cannot be bitwise-identical, i.e. it is the\nprice of the determinism guarantee — fucina reads 44% fewer weight bytes and both\nengines run ~2× above the memory floor there. **We keep the guarantee.**\n\n---\n\n## 🖥️ Hardware support\n\nfucina is built and tested for **exactly one accelerator**: the **NVIDIA DGX Spark GB10**.\n\n| | |\n|---|---|\n| **Supported GPU** | NVIDIA **DGX Spark GB10** (Grace-Blackwell) |\n| **Compute capability** | **`sm_121`**, compiled as **`sm_121a`** (arch-specific GB10 features: FP8 / NVFP4 block-scaled MMA, `tcgen05`) |\n| **Memory** | GB10's **128 GB unified LPDDR5X** (CPU+GPU shared) |\n| **CUDA** | **13.0** at `/usr/local/cuda-13` |\n\n**Other GPUs are not supported as built.** The Makefile pins `CUDA_ARCH := sm_121a` and the fast\npaths depend on GB10-class tensor-core features:\n\n- **Consumer Blackwell** (RTX 50-series, `sm_120`) and **datacenter Blackwell** (B100/B200,\n `sm_100`) are **not targeted or tested** — they would at minimum need a `CUDA_ARCH` change and\n FP8/NVFP4 re-validation. Treat as unverified.\n- **Pre-Blackwell GPUs** (Hopper, Ada, Ampere, …) lack the FP8/NVFP4 block-scaled MMA paths the\n engine relies on and are out of scope.\n\n> [!TIP]\n> **Context vs. memory (Gemma-4).** The sliding-window KV cache (40 of 48 layers) is a capped\n> **ring buffer**, so it does **not** grow with `--ctx`: ~**1.5 GiB** at the default ring size\n> (8192 slots), tunable via `FUCINA_SLIDING_RING`. Only the 8 global-attention layers' KV scales\n> with context (~2.0 GiB at the full 262144). Total FP8 KV cache (1 B/element, K+V): **~3.5 GiB at\n> 262144**, **~2.5 GiB at 131072** — down from ~54 GiB before the ring. To shrink the *weights*\n> footprint, set `FUCINA_NO_PACKED=1` to drop the repacked-Q4_0 decode copy (**−~7 GiB**, ~2–3%\n> slower decode). The engine clamps `--ctx` to 262144. See [Environment toggles](#-http-api).\n\n> [!TIP]\n> **Memory sizing (Qwen).** There is no hardcoded minimum-VRAM figure for the 35B-A3B MoE\n> checkpoints. `--gpu-mem-util` (default `0.90`) caps total device memory the engine will use as a\n> fraction of the GPU's total; the paged-KV pool is auto-sized to whatever free memory remains\n> after weights, and `--ctx` / the paged pool auto-cap to fit. If a Qwen checkpoint fails to start\n> with a paged-KV allocation error, lower `--ctx` or raise `--gpu-mem-util` first before assuming\n> the checkpoint doesn't fit at all. The 27B dense and 35B-A3B MoE FP8 checkpoints, plus the\n> smaller NVFP4 ModelOpt build, are all default-Q4_K-requanted for the mixer/FFN/expert weights at\n> load (see [Models](#-models)), which shrinks the resident set relative to their on-disk FP8 size.\n\n### Requirements\n\n- **NVIDIA DGX Spark GB10** — Blackwell, compute capability `sm_121`.\n- **CUDA 13.0** at `/usr/local/cuda-13`.\n- **Go 1.26** (the Makefile expects `/usr/local/go/bin/go`).\n- A supported checkpoint ([download below](#-models)): **Gemma 4 12B** GGUF (**Q4_0 QAT** or\n **Q8_0**) or NVFP4 safetensors; or a **Qwen3 / Qwen3.5 / Qwen3.6** GGUF, official FP8-block\n safetensors, or NVIDIA ModelOpt NVFP4/FP8 checkpoint.\n- *(Optional, Gemma-4 only)* the **Gemma 4 MTP draft head** GGUF for\n [speculative decoding](#-speculative-decoding).\n\n---\n\n## 📥 Models\n\nfucina loads **local model files** (GGUF or safetensors) and has **no model-download logic of its\nown** — fetch weights with `hf` (or `huggingface-cli download …`, equivalent) and pass the path\nwith `-m`. The architecture and quant format are **auto-detected from the file itself** — there is\nno `--model-type`/`--arch` flag. *Repo ids and filenames below are verified against Hugging Face.*\n\n```sh\npip install -U \"huggingface_hub[cli]\" # provides the `hf` command\n```\n\n### Model support matrix\n\n| Family | `-m` accepts | Detected via | Batching |\n|---|---|---|---|\n| Gemma 4 12B | Q4_0-QAT / Q8_0 GGUF; NVFP4 safetensors | GGUF falls through to Gemma when `general.architecture` isn't a `qwen*` key; NVFP4 dir/`config.json` | opt-in (`--batch`) |\n| Qwen3 dense | GGUF | `general.architecture = \"qwen3\"` | mandatory (auto) |\n| Qwen3 MoE | GGUF | `general.architecture = \"qwen3moe\"` | mandatory (auto) |\n| Qwen3.5 / Qwen3.6 hybrid, dense | GGUF, or official FP8-block safetensors | GGUF `\"qwen35\"`; safetensors `config.json` `model_type` contains `\"qwen3_5\"` + `quant_method` `fp8`/`modelopt`, no expert tensors | mandatory (auto) |\n| Qwen3.5 / Qwen3.6 hybrid, MoE (A3B) | official FP8-block safetensors, or NVIDIA ModelOpt mixed NVFP4/FP8 safetensors | same as above, plus an `mlp.experts.0.gate_proj.weight` tensor | mandatory (auto) |\n| DiffusionGemma 26B-A4B | Q4_K_M GGUF (via `-dm`) | separate engine, explicit flag | N/A (block diffusion) |\n\n**\"Qwen3.6\" is a checkpoint generation, not a separate code path** — `Qwen/Qwen3.6-27B-FP8`,\n`Qwen/Qwen3.6-35B-A3B-FP8`, and `nvidia/Qwen3.6-35B-A3B-NVFP4` are all recognized and served through\nthe same Qwen3.5 hybrid detector and loader (identical `qwen3_5` architecture).\n\n**Continuous batching is mandatory for every Qwen3/3.5/3.6 checkpoint** (no single-flight path\nexists for Qwen); fucina detects this and turns `--batch`/`--paged-kv` on for you automatically — a\nbare `fucina -m ` just works. See\n[Continuous batching & paged KV](#-continuous-batching--paged-kv).\n\n### HuggingFace checkpoint paths — the robust way\n\nFor any safetensors checkpoint (Qwen or Gemma NVFP4), **always download with `--local-dir` and\npoint `-m` at that directory**:\n\n```sh\nhf download --local-dir ./models/\nfucina -m ./models/ ...\n```\n\nThis gives you a flat directory with `config.json` + weight shards + `tokenizer.json` at the top\nlevel, so both the weight loader and the tokenizer auto-discovery (below) find everything with no\n`--tokenizer` flag and no ambiguity.\n\nIf you instead used the bare `hf download ` cache (no `--local-dir`), the files land under\n`~/.cache/huggingface/hub/models----/snapshots//` — pass **that snapshot\ndirectory** to `-m`, not the `models--...` parent. The Qwen3.5/3.6 FP8/ModelOpt loader specifically\nresolves a bare repo-root parent too (it globs `snapshots/*/model.safetensors.index.json`), but the\nGo-side tokenizer auto-discovery does **not** do that glob and will fail to find `tokenizer.json`\nif you point `-m` at the parent — you'd need an explicit `--tokenizer /tokenizer.json`\nin that case. The generic Gemma-4 NVFP4 loader has no repo-root resolution at all and needs the\nreal snapshot-shaped directory. Using `--local-dir` sidesteps all of this.\n\n**Tokenizer auto-discovery, in short:** GGUF models never need `--tokenizer` (vocab is inline).\nSafetensors/NVFP4/FP8 checkpoints need a `tokenizer.json`, found automatically as a sibling of `-m`\n(per the path rule above) or passed explicitly via `--tokenizer `.\n\n**1 · Dense Gemma 4 12B** *(required)* — Q4_0 **QAT** is the recommended path (the official\nquantization-aware-trained weights: Q4_0 layers with a Q6_K tied LM head). Q8_0 also works.\n\n```sh\n# Q4_0 (QAT) — recommended (official Google QAT GGUF)\nhf download google/gemma-4-12B-it-qat-q4_0-gguf \\\n gemma-4-12b-it-qat-q4_0.gguf --local-dir ./models\n\n# or Q8_0 (unsloth)\nhf download unsloth/gemma-4-12b-it-GGUF \\\n gemma-4-12b-it-Q8_0.gguf --local-dir ./models\n```\n\n**2 · MTP draft head** *(optional, for [speculation](#-speculative-decoding))* — a small\n(~444 MB, **Q8_0**) separate GGUF. It is a GGUF build of Google's official assistant\n**`google/gemma-4-12B-it-assistant`** (architecture `gemma4-assistant`, 4 layers), shipped inside\nunsloth's GGUF repo under `MTP/`.\n\n```sh\nhf download unsloth/gemma-4-12b-it-GGUF \\\n MTP/gemma-4-12b-it-Q8_0-MTP.gguf --local-dir ./models\n# → ./models/MTP/gemma-4-12b-it-Q8_0-MTP.gguf\n```\n\n**3 · DiffusionGemma 26B-A4B** *(optional, for the [`-dm` engine](#-diffusiongemma))*\n\n```sh\nhf download unsloth/diffusiongemma-26B-A4B-it-GGUF \\\n diffusiongemma-26B-A4B-it-Q4_K_M.gguf --local-dir ./models\n```\n\n> [!NOTE]\n> Google's `google/gemma-4-12B-it-assistant` is the canonical MTP head but is published as\n> **safetensors**; fucina's `--assistant` needs a **GGUF** build (the unsloth file above).\n> If you place a dense model at `./gemma-4-12b-it.gguf`, `./model.gguf`, or `./gguf/model.gguf`,\n> fucina finds it automatically when `-m` is omitted. The reported model id is derived from the\n> GGUF filename.\n\n### Gemma 4 E4B (Per-Layer Embeddings)\n\nfucina also runs **Gemma-4-E4B-it** — the effective-4B variant (Per-Layer Embeddings, KV-sharing,\nruntime-detected dims) — through a separate engine, auto-selected from the GGUF metadata (no flag).\nDecode reads the original **Q4_0 QAT nibbles** directly via the shared dp4a MMVQ kernels and a\nnative Q6_K tied head; the BF16 projections are freed after load, so the weight footprint is\n**~7 GB** (see [Performance](#-performance)). Q4_K_M / Q5_K / BF16 GGUFs also load.\n\n```sh\n# Q4_0 (QAT) — official Google E4B GGUF\nhf download google/gemma-4-E4B-it-qat-q4_0-gguf \\\n gemma-4-E4B_q4_0-it.gguf --local-dir ./models\n\n./fucina -m ./models/gemma-4-E4B_q4_0-it.gguf --interactive # same REPL commands as the 12B\n```\n\nThe E4B REPL shares the dense command surface (`/help`, `/thinking`, `/stats`, `/reset`, `/quit`)\nand `--thinking off|on|low|medium|high|xhigh`. Server mode and the MTP draft head are not wired for\nE4B yet.\n\n### Native NVFP4 (safetensors) — Gemma 4\n\nfucina also runs a **natively NVFP4-quantized** Gemma 4 loaded straight from a HuggingFace\nsafetensors checkpoint (e.g. `RedHatAI/gemma-4-12B-it-NVFP4`) — a single FP4 weight store feeding\nboth the cuBLASLt block-scaled prefill and a fused FP4 decode GEMV, with the checkpoint's own\n`tokenizer.json` read natively (no `--tokenizer`). Point `-m` at the checkpoint directory (see\n[HuggingFace checkpoint paths](#huggingface-checkpoint-paths--the-robust-way) above — use\n`--local-dir` so this \"just the directory\" usage works without a snapshot-path gotcha):\n\n```bash\n./fucina -m /path/to/gemma-4-12B-it-NVFP4 -p \"Explain photosynthesis.\" -n 128\n```\n\nBoth compressed-tensors (`RedHatAI/*`) and ModelOpt (`nvidia/*-FP4`) naming are auto-detected. See\n[docs/nvfp4-safetensors.md](docs/nvfp4-safetensors.md) for the schema, architecture, performance,\nand limits. **This is a different loader from the Qwen3.5/3.6 NVFP4 path below** — this one is\nGemma-4-only, single-expert (no MoE), and has no repo-root resolution (needs the real snapshot dir).\n\n> [!TIP]\n> **Speculative decoding with NVFP4.** The MTP draft head (`--assistant`) accepts NVFP4's tokens at\n> a lower rate (~42%) than the original model (~89%): the assistant is matched to the original\n> weights, not the NVFP4 checkpoint. For NVFP4, **prompt-lookup speculation** (model-agnostic,\n> enabled by default) is the recommended speculative path. MTP still works, but the throughput gain\n> is smaller (~28 tok/s vs ~57 with Q4_0+MTP). See\n> [docs/nvfp4-safetensors.md](docs/nvfp4-safetensors.md) for details.\n\n### 4 · Qwen3 / Qwen3.5 / Qwen3.6 (dense and MoE)\n\nAll Qwen checkpoints below are served through the same binary with **no model-select flag** —\njust point `-m` at the file or directory. Continuous batching auto-enables itself (see\n[Continuous batching & paged KV](#-continuous-batching--paged-kv)). For the full detection\nreference, the official-FP8-vs-ModelOpt distinction, chat-dialect/tool-calling details, and a\ndiagnostics table, see [docs/qwen-models.md](docs/qwen-models.md).\n\n**Qwen3 / Qwen3.5 GGUF (dense or MoE)** — any standard GGUF conversion works: the loader reads\n`general.architecture` (`qwen3`, `qwen3moe`, or `qwen35`) directly from the file, and no\n`--tokenizer` is needed (vocab is inline in the GGUF, same as Gemma).\n\n```bash\n./fucina -m ./models/.gguf --host 0.0.0.0 --port 8080\n```\n\n**Qwen3.5 / Qwen3.6, official FP8-block safetensors** (DeepSeek-V3-style block quant) — dense or\nMoE, distinguished automatically by the presence of expert tensors:\n\n```sh\n# Qwen3.5-9B (dense) / Qwen3.5-35B-A3B (MoE)\nhf download Qwen/Qwen3.5-9B-FP8 --local-dir ./models/Qwen3.5-9B-FP8\nhf download Qwen/Qwen3.5-35B-A3B-FP8 --local-dir ./models/Qwen3.5-35B-A3B-FP8\n\n# Qwen3.6-27B (dense) / Qwen3.6-35B-A3B (MoE) — same architecture/loader as Qwen3.5 above\nhf download Qwen/Qwen3.6-27B-FP8 --local-dir ./models/Qwen3.6-27B-FP8\nhf download Qwen/Qwen3.6-35B-A3B-FP8 --local-dir ./models/Qwen3.6-35B-A3B-FP8\n\nfucina -m ./models/Qwen3.6-35B-A3B-FP8 --host 0.0.0.0 --port 8080\n```\n\n**Qwen3.6-35B-A3B, NVIDIA ModelOpt NVFP4/FP8 (MIXED_PRECISION)** — this is a **different\ncheckpoint and loader** from the official FP8 build above: NVIDIA's ModelOpt re-quantizes the\nMoE experts, shared expert, and LM head to native **NVFP4**, while keeping attention and the\nGated-DeltaNet mixer in **per-tensor FP8** — smaller on-disk and resident footprint than the\nofficial all-FP8 checkpoint, at a small accuracy cost from the 4-bit experts. Detected the same\nway (`config.json` `model_type`/`quant_method`), with the FP8-vs-NVFP4 choice made **per tensor**\nfrom which scale sibling is present — no flag selects it.\n\n```sh\nhf download nvidia/Qwen3.6-35B-A3B-NVFP4 --local-dir ./models/Qwen3.6-35B-A3B-NVFP4\nfucina -m ./models/Qwen3.6-35B-A3B-NVFP4 --host 0.0.0.0 --port 8080\n```\n\n> [!NOTE]\n> **Perf default:** dense **and** MoE Qwen3.5/3.6 checkpoints get their attention/GDN mixer and\n> FFN/expert weights requanted to **Q4_K** at load time (smaller resident weight set, faster\n> decode than the on-disk FP8). Opt back into pure FP8 with `FUCINA_MOE_FP8=1` if you need to\n> compare against the unquantized checkpoint.\n\n> [!NOTE]\n> **Tool calling / thinking / structured output** work over the standard OpenAI wire shape for\n> Qwen too — see [HTTP API](#-http-api). `response_format` (`json_object` and the supported\n> `json_schema` subset) carries an independent grammar state per continuous-batch slot. See\n> [batched structured output](docs/batched-structured-output.md) for correctness and GB10 evidence.\n\n---\n\n## 🧵 Continuous batching & paged KV\n\nIndependent requests are served through a per-step scheduler over a **paged, multi-sequence KV\ncache** (vLLM-style block table + free-list) instead of one lock held for the whole request —\nconcurrent clients share each batched forward pass, and CUDA graphs are captured per active batch\nsize.\n\n- **Gemma-4: opt-in.** Off by default (the default single-flight path is faster for one client at\n a time and has MTP speculation, which the batch path lacks). Turn it on with `--batch` (implies\n `--paged-kv`) — equivalent to the legacy `FUCINA_PAGED_KV=1 FUCINA_BATCH=1` env pair — when you\n are actually serving concurrent clients.\n- **Qwen3 / Qwen3.5 / Qwen3.6: mandatory, auto-enabled.** There is no single-flight path for Qwen\n (its prefill entry points decline), so fucina detects the architecture and turns batching on for\n you — `fucina -m ` just works, with or without `--batch`. If the paged-KV pool\n fails to allocate, startup fails fast with a message telling you to lower `--ctx` or raise\n `--gpu-mem-util`, rather than serving requests that would 500.\n- **Speculative decoding inside the batch path** is prompt-lookup only (n-gram, free) and applies\n automatically to **dense** checkpoints (Gemma or Qwen); it's disabled for **MoE** checkpoints.\n The Gemma-4 MTP draft head (`--assistant`) is not exercised in the batch path.\n- **Burst-admission coalescing:** when the scheduler wakes from idle, it holds a short escalating\n window (a few ms, capped at 150 ms) so near-simultaneous requests land in the same batch instead\n of admitting one-by-one — unconditional scheduler behavior, no flag.\n- **Structured output:** `response_format`/`json_schema` works in the batch path. Constrained rows\n use exact one-token decode plus per-slot host grammar masking; speculation is disabled only for\n shared steps containing a constrained row. Adapters without exact logits still fail closed.\n\nSee [`docs/continuous-batching.md`](docs/continuous-batching.md) for the full design (paged-KV\nallocator, split-K paged attention, per-batch-size CUDA graphs) and\n[`docs/qwen35-beat-vllm-plan.md`](docs/qwen35-beat-vllm-plan.md) for the Qwen-specific measurement\nrecord.\n\n---\n\n## 🔗 Experimental Phase-E layer sharding\n\nQwen3.5 has an exact fp32 residual-stream boundary for layer ranges and a persistent\n`FCNDIST1` TCP coordinator/worker path. The current CLI supports **one-shot,\ntoken-sequential** inference only; it rejects multi-token frames rather than violating GDN\nrecurrence order. See [`docs/phase-e-distributed.md`](docs/phase-e-distributed.md) for commands,\nprotocol/lifecycle guarantees, and the hardware gates.\n\nThis is **not yet a multi-node capacity claim**: every process still loads the full checkpoint,\nthe HTTP scheduler is not distributed, and physical two-GB10 correctness/performance gates have\nnot run. Range-filtered weight residency is the blocker for models that cannot fit one node.\n\n---\n\n## 🎯 Speculative decoding\n\nSpeculative decoding is **on by default** (`--spec`). Each step *drafts* candidate tokens cheaply,\nthen the model verifies them all in **one batched weight pass** and keeps the longest matching\nprefix — so accepted drafts cost far less than one decode each. Every verified position is drawn\nfrom the **target model's own distribution**, so the output is **identical to plain decoding** at\nthe same settings — drafting only changes speed, never the result.\n\nThis section describes the Gemma-4 single-flight path, where both drafters below are available.\nQwen checkpoints run through the continuous-batching path instead (see\n[Continuous batching & paged KV](#-continuous-batching--paged-kv)): prompt-lookup speculation\napplies automatically to **dense** Qwen checkpoints, is off for **MoE**, and the MTP draft head\nbelow (`--assistant`) does not apply to Qwen at all.\n\nTwo drafters work together (Gemma-4):\n\n| Drafter | Needs | Best at |\n|---------|-------|---------|\n| **Prompt-lookup** | nothing (on with `--spec`) | repetitive / structured / code-like text — finds where recent context recurs earlier (consensus n-gram match). Free, host-side. |\n| **MTP draft head** | `--assistant .gguf` | **novel** text — a 4-layer multi-token-prediction head (Google's `gemma-4-12B-it-assistant`, ~444 MB Q8_0) running over the shared frozen KV cache. llama.cpp's `--spec-type draft-mtp` equivalent. |\n\nWhen both are available the engine prefers MTP for novel text and only lets a strong prompt-lookup\ndraft displace it; draft length adapts per step from each drafter's running acceptance rate\n(clamped to `[2, --draft-k]`).\n\n```sh\n# Enable MTP (server, one-shot, or REPL — just add --assistant)\nfucina -m ./models/gemma-4-12b-it-qat-q4_0.gguf \\\n --assistant ./models/MTP/gemma-4-12b-it-Q8_0-MTP.gguf \\\n --ctx 32768 --host 0.0.0.0 --port 8080\n```\n\n- **Works at any temperature** — correct for both greedy (`--temp 0`) and sampling (`--temp > 0`).\n- **Observe it live** — `/metrics` → `speculation` (`accept_rate`, `tokens_per_forward`); CLI runs\n print `[mtp]` / `[lookup]` stats at the end.\n- **Falls back** to per-token decode for requests with text `stop` strings (host-side trimming).\n- **Disable:** `--spec=false`. Graph escape hatches: `FUCINA_NO_DECODE_GRAPH=1`,\n `FUCINA_NO_BATCHED_GRAPH=1`. The draft head adds ~0.4 GB VRAM.\n\n> The `--spec` flag's built-in help still reads *\"greedy/temp=0 only\"* — that wording is stale; the\n> engine fully supports sampling, including on-GPU repeat-penalty at `--temp > 0`.\n\n---\n\n## 🌐 HTTP API\n\n| Endpoint | Description |\n|----------|-------------|\n| `POST /v1/chat/completions` | OpenAI chat completions — streaming + non-streaming, tool calls, thinking channel |\n| `POST /v1/messages` | Anthropic Messages API — streaming + non-streaming, tool use/results, thinking blocks |\n| `POST /v1/completions` | Legacy raw-prompt completions (handled by the chat path) |\n| `GET /v1/models` | Lists the loaded model id |\n| `POST /v1/embeddings` | Stub — returns an empty data list |\n| `GET /health`, `/healthz` | Liveness + KV-cache stats (hits, misses, hit rate, cached tokens) |\n| `GET /readyz` | Readiness — checks the tokenizer + engine are loaded; `503` when not serviceable |\n| `GET /metrics` | KV/context utilization, prefix-cache hit rate, prefill/decode throughput, `speculation`, `requests_detail` (total, errors, avg latency, avg TTFT), `saturation` (in-flight / max) |\n\n`/v1/*` routes accept an optional `Authorization: Bearer ` (see `--api-key`); the Anthropic\nendpoint also accepts `x-api-key: `. `/health`, `/healthz`, `/readyz`, and `/metrics` are always\nopen. Every response carries an `X-Request-Id` (echoed from the request when present) for log\ncorrelation.\n\nThe chat dialect (Gemma-4 native format, or Qwen ChatML + Qwen3-Coder XML tool calls) is picked\n**automatically from the loaded vocab** — no flag. The wire format is OpenAI-shaped either way:\n\n```sh\ncurl http://localhost:8080/v1/chat/completions -d '{\n \"messages\": [{\"role\": \"user\", \"content\": \"weather in Paris?\"}],\n \"tools\": [{\n \"type\": \"function\",\n \"function\": {\n \"name\": \"get_weather\",\n \"description\": \"Get the weather\",\n \"parameters\": {\n \"type\": \"object\",\n \"properties\": {\"city\": {\"type\": \"string\"}},\n \"required\": [\"city\"]\n }\n }\n }],\n \"max_tokens\": 64\n}'\n```\n\nThe response carries standard `choices[0].message.tool_calls[].function.{name,arguments}` and\n`finish_reason: \"tool_calls\"` regardless of dialect. For Qwen, the server auto-detects Qwen3.6's\nXML-shaped output (`V`)\nor the legacy Qwen3/3.5 JSON body inside `` and maps either to that same shape. Force a\nspecific tool with\n`\"tool_choice\": {\"type\":\"function\",\"function\":{\"name\":\"get_weather\"}}`.\n\nReasoning/thinking works the same for both dialects: `\"reasoning_effort\": \"low\"|\"medium\"|\"high\"` or\n`\"thinking\": true/false` in an OpenAI request; the model's reasoning appears in\n`choices[0].message.reasoning_content`.\n\nThe same server can be used directly with Anthropic SDKs by setting their base URL to the fucina\nserver and calling `/v1/messages`. Anthropic `system`, content blocks, `tools`, `tool_choice`,\n`stop_sequences`, `thinking`, and SSE event shapes are translated to the same generation\npath. For example:\n\n```sh\ncurl http://localhost:8080/v1/messages \\\n -H 'anthropic-version: 2023-06-01' \\\n -H 'content-type: application/json' \\\n -d '{\n \"model\": \"local\",\n \"max_tokens\": 64,\n \"messages\": [{\"role\": \"user\", \"content\": \"Hello\"}],\n \"stream\": true\n }'\n```\n\nFor restart-safe conversations, start with `--session-dir DIR` and include\n`\"session\": \"NAME\"` in chat or legacy completion requests. Gemma saves flat KV;\nQwen3.5/3.6 saves each batch slot's full-attention KV plus GDN recurrent/conv state.\nThe next request must render a strict extension of the saved token history. See\n[`docs/session-persistence.md`](docs/session-persistence.md) for format and safety details.\n\n> [!IMPORTANT]\n> `response_format`/`json_schema` uses host-side grammar masking. In continuous batching each slot\n> owns its own FSM and RNG; constrained shared steps use the exact host-input ABI (no speculative\n> drafts or Qwen GPU token splice). This is correctness-first and copies full logits to the host,\n> so it is slower than ordinary on-device sampling. Engines without that ABI retain the fail-closed\n> `501 unsupported_under_batching` guard.\n\n---\n\n## 🏗️ Architecture\n\n```\n┌───────────────────────────────────────────────────────────────────┐\n│ cmd/fucina CLI: server / one-shot / interactive │\n├───────────────────────────────────────────────────────────────────┤\n│ internal/server OpenAI/Anthropic HTTP APIs + KV cache │\n│ internal/server/batch continuous-batching scheduler (Qwen; opt-in │\n│ for Gemma-4) │\n│ internal/chat chat-template + tool-call dialects (Gemma/Qwen)│\n│ internal/tokenizer SentencePiece/Unigram (Gemma) + native HF BPE │\n│ (Qwen tokenizer.json) │\n├───────────────────────────────────────────────────────────────────┤\n│ internal/engine/cuda CGO bridge (Go ⇄ C) │\n├───────────────────────────────────────────────────────────────────┤\n│ cuda/ CUDA C++ kernels: Gemma-4, Qwen3/3.5/3.6, │\n│ GDN mixer, MoE grouped-expert GEMM (sm_121a) │\n└───────────────────────────────────────────────────────────────────┘\n```\n\n- **`cuda/`** — the forward pass for every architecture: quantized GEMV/GEMM (Q4_0/Q6_K/Q8_0/Q4_K\n `dp4a`), FP8/NVFP4 flash attention, the Qwen3.5/3.6 Gated-DeltaNet mixer, MoE grouped-expert GEMM,\n RoPE, RMSNorm, GPU sampling, CUDA-graph capture/replay, and paged multi-sequence batching.\n Compiled to `libfucina.a`. Key files: `gemma4_detect.h` (arch/format detection),\n `qwen35_fp8_loader.h` (Qwen3.5/3.6 FP8-block + ModelOpt NVFP4 safetensors loader),\n `nvfp4_loader.h`/`safetensors.h` (Gemma-4 NVFP4), `gemma4_kernels.cu` (everything else).\n- **`internal/engine/cuda`** — the CGO bridge, wrapping the opaque CUDA engine as a Go `Engine`\n (`NewEngine`, `Prefill`, `Decode`, `GenerateSpec`, `LoadAssistant`, `IsQwen3Family`,\n `SeqAdd`/`StepBatch`, …).\n- **`internal/tokenizer`** — Gemma SentencePiece/Unigram (vocab 262144, from the GGUF) and a native\n HuggingFace BPE reader (`tokenizer.json`, byte-fallback) for Qwen/NVFP4/FP8 checkpoints.\n- **`internal/chat`** — the `Dialect` interface (`Gemma`, `Qwen`), auto-selected from the loaded\n vocab; chat-template rendering and tool-call parsing/formatting for each.\n- **`internal/server`** — the HTTP server, thinking channel, `response_format`/JSON-schema\n constrained decoding (`internal/grammar`), and prefix-reuse KV cache (single-flight path).\n- **`internal/server/batch`** — the continuous-batching scheduler: admission/burst-coalescing,\n per-step batch build, per-sequence spec gating (dense-only), eviction.\n- **`cmd/fucina`** — CLI entry point and the three run modes.\n\n> [!IMPORTANT]\n> cgo does **not** hash the contents of the `-lfucina` static archive, so a plain `go build` can\n> silently relink a stale binary against an updated `libfucina.a`. The Makefile defends against this:\n> it removes the old binary, rebuilds with `go build -a`, and asserts the device-upload path is\n> present in the binary.\n\n---\n\n## ⚙️ Configuration\n\n\nImportant flags (run fucina --help for the full list)
\n\n| Flag | Default | Description |\n|------|---------|-------------|\n| `-m, --model` | (required) | GGUF file, or a safetensors file/directory (FP8-block, NVFP4, or ModelOpt mixed); architecture and quant auto-detected |\n| `--tokenizer` | (auto) | `tokenizer.json` or `.gguf` to source vocab from; only needed if auto-discovery (sibling of `-m`) fails |\n| `--ctx` | `262144` | Context size in tokens (max/default 262144; lower it to save memory) |\n| `--temp` | `1.0` | Sampling temperature (gemma-4 default; `0` = greedy) |\n| `--top-k` | `64` | Top-K sampling (gemma-4 default) |\n| `--top-p` | `0.95` | Top-P / nucleus sampling (gemma-4 default) |\n| `--thinking` | `off` | Default reasoning channel: `off`/`on`/`low`/`mid`/`high`/`xhigh` (Gemma and Qwen) |\n| `--assistant` | (none) | Gemma-4 MTP draft-head GGUF for speculation (Gemma-4 only; no effect on Qwen) |\n| `--spec` | `true` | Speculative decoding (prompt-lookup; MTP too when `--assistant` is set, Gemma-4 single-flight only) |\n| `--draft-k` | `6` | Max speculative draft length per step |\n| `--paged-kv` | `false` | Allocate the paged multi-sequence KV pools; auto-forced on for any Qwen3 checkpoint |\n| `--batch` | `false` | Continuous batching over the paged engine (implies `--paged-kv`); auto-forced on for any Qwen3 checkpoint, opt-in for Gemma-4 |\n| `--dist-listen` | (none) | Experimental Phase-E headless worker listen address; requires `--dist-layers` |\n| `--dist-workers` | (none) | Experimental one-shot coordinator: comma-separated worker addresses in layer order |\n| `--dist-layers` | (none) | This process's strict `lo:hi` Qwen3.5 layer range |\n| `--dist-final` | `false` | Mark a worker as the final range; applies output norm/head and returns logits |\n| `--gpu-mem-util` | `0.90` | Fraction of total GPU memory the engine may use; caps `--ctx`/paged-pool sizing to fit |\n| `--host` | `127.0.0.1` | Server listen address |\n| `--port` | `8080` | Server port |\n| `--api-key` | (none) | Bearer token required on `/v1/*` (constant-time; reads `FUCINA_API_KEY` if unset). Empty = auth off (localhost dev) |\n| `--max-concurrent` | `4` | Admission-queue depth (in-flight + waiting); excess requests get `503` |\n| `--max-output-tokens` | `0` | Absolute per-request output-token ceiling (independent of context window); `0` = no extra cap |\n| `--session-dir` | (none) | Directory for named disk sessions supplied as `\"session\": \"NAME\"`; supports Gemma single-flight and Qwen3.5/3.6 batching |\n\n \n\n\nEnvironment toggles
\n\n| Variable | Effect |\n|----------|--------|\n| `FUCINA_NO_DECODE_GRAPH=1` | Disable CUDA-graph capture for single-token decode |\n| `FUCINA_NO_BATCHED_GRAPH=1` | Disable CUDA-graph capture for the K-row batched verify |\n| `FUCINA_NO_PACKED=1` | Drop the repacked-Q4_0 decode-GEMV weight copy (Gemma-4). **Frees ~7 GiB VRAM** at the cost of ~2–3% slower decode; output is bit-identical. The repacked copy is a second, coalesced-load layout of the Q4_0 projection weights kept resident only to speed the bandwidth-bound decode hot path. Recommended on memory-constrained hosts. |\n| `FUCINA_SLIDING_RING=N` | Sliding-window KV ring capacity in tokens (default **8192**, floored at `window+spec_max`), Gemma-4 only. Caps the sliding cache so it stays ctx-independent (~1.5 GiB at 8192). `N` also bounds how far a prefix-reuse rewind stays exact — deeper rewinds (e.g. editing context older than `N-1024` tokens) fall back to a full re-prefill. With `--ctx ≤ N` behavior is identical to the old flat cache. Lower = less VRAM; higher = deeper exact rewind (~+190 MiB per +1024). |\n| `FUCINA_PAGED_KV=1` | Allocate the paged KV pools (block table + free-list), capacity-sized to free VRAM. Prerequisite for continuous batching; equivalent to `--paged-kv`. Always on for Qwen3. |\n| `FUCINA_BATCH=1` | Route the server through the continuous-batching scheduler (needs `FUCINA_PAGED_KV=1`): independent sequences share each batched forward instead of the per-request lock. Equivalent to `--batch`; always on for Qwen3. |\n| `FUCINA_MOE_FP8=1` | Disable the default Q4_K requant of the Qwen3.5/3.6 attention/GDN mixer and FFN/expert weights; serve pure FP8 instead. |\n| `FUCINA_QWEN35_FP4=1` | Opt-in NVFP4 activations for the Qwen3.5/3.6 hybrid path (in addition to the default weight quant). |\n| `FUCINA_NO_BATCH_SPEC=1` | Disable prompt-lookup speculation inside the continuous-batching scheduler (Gemma or Qwen dense). |\n| `FUCINA_PAGED_MAXCTX=N` | Override the paged-KV pool sizing context (default 32768, clamped to `--ctx`). |\n| `FUCINA_PAGED_MAXSEQS=N` | Override the auto-computed max concurrent sequences the paged pool backs (auto range 1–64). |\n| `FUCINA_KV_NVFP4=1` | Quantize the KV cache to **NVFP4** precision (native Blackwell FP4 E2M1 + per-16 E4M3 block scale). Memory-only (~1.78× smaller KV), small accuracy cost; default OFF keeps flat FP8 byte-identical. |\n| `FUCINA_NO_WARMUP_PASS=1` | Skip the one-time startup warmup pass |\n| `FUCINA_DEBUG=1` | Dump request bodies + rendered prompts to `/tmp/fucina_debug.log` |\n\n \n\n---\n\n## 🌫️ DiffusionGemma\n\n\nRunning the 26B-A4B text-diffusion MoE model
\n\nDiffusionGemma is loaded with **`-dm`** (`--diffusion-model`) instead of `-m`. That routes to a\nseparate diffusion CUDA engine and enables the **NVFP4 Tensor-Core MoE experts**. All three run\nmodes and the OpenAI API work exactly as for the dense model.\n\n```sh\n# Server\nfucina -dm ./models/diffusiongemma-26B-A4B-it-Q4_K_M.gguf --ctx 8192 --host 0.0.0.0 --port 8080\n\n# One-shot\nfucina -dm ./models/diffusiongemma-26B-A4B-it-Q4_K_M.gguf -p \"Write a haiku about the ocean.\"\n\n# Faster generation, lower quality: fewer denoise steps per block (default 48)\nfucina -dm ./models/diffusiongemma-26B-A4B-it-Q4_K_M.gguf --denoise-steps 16 -p \"Explain hashing.\"\n```\n\n**How it differs from the dense model:**\n\n- **Block diffusion, not autoregressive.** Output is generated by iteratively *denoising* a fixed\n **256-token canvas**, then chaining blocks until an end-of-turn/EOS or `max_tokens`.\n- **No token streaming.** A whole block is denoised at once (REPL shows `denoising…`; the server\n emits the block as one SSE delta).\n- **Two throughput figures** are reported: `tokens_per_second` (delivered) and\n `canvas_tokens_per_second` (raw denoising rate). Delivered is the apples-to-apples number.\n- **Context auto-caps to GPU memory** — pass a large `--ctx`; the engine logs the capped `max_prompt`.\n\n| Flag | Default | Description |\n|------|---------|-------------|\n| `-dm, --diffusion-model` | (required) | DiffusionGemma GGUF; routes to the diffusion engine + NVFP4 MoE |\n| `--denoise-steps` | `48` | Denoise steps per block; **lower = faster, lower quality** (blocks converge ~13–16) |\n| `--fp4-moe` | (implied) | NVFP4 MoE experts (on automatically with `-dm`; also usable with `-m`) |\n| `--ctx` | `8192` | Context ceiling (auto-capped to free GPU memory) |\n\n \n\n---\n\n## 🧪 Testing\n\n```sh\n# CPU-only unit tests (no GPU, no cgo, no model file required) — this is what CI runs\ngo test ./internal/server/ ./internal/server/batch/ ./internal/tokenizer/ ./internal/sampler/ ./internal/chat/\nmake check # go vet + the same pure-Go tests + gofmt check\n\n# GPU smoke test (requires the DGX Spark GB10)\nmake smoke # builds, then: fucina --prompt \"Hello, world!\" --predict 32 --temp 0\n\n# GPU correctness + performance (requires the GB10 + a Gemma model; MODEL=… overridable)\nmake bench # batch==single self-test + greedy byte-identity, then prefill/decode tok/s\nmake paged-kv-device-test # paged KV reads bit-identical to the contiguous cache\nmake packed-kv-test # packed 4.5-bit NVFP4 KV storage bit-identical to the fake-quant\nmake kv-quant-explore # offline FP8 / NVFP4 / TurboQuant codec comparison (host-only)\n\n# Qwen3/3.5/3.6 GPU correctness gates (require the GB10 AND real downloaded checkpoints at\n# hardcoded /opt/spark/models/... paths — override with QWEN3_DENSE_MODEL=/QWEN3MOE_MODEL=/\n# QWEN35_MODEL=/QWEN35_FP8_MODEL=/QWEN35_MOE_FP8_MODEL=. NOT runnable in CI or without hardware+weights.\nmake gpu-gates # Qwen3 dense/MoE parity + spec + prefix/suffix regression gates\nmake qwen35-batch-test # Qwen3.5/3.6 paged-batch + CUDA-graph decode gate\nmake qwen35-burst-test # diverse-prompt burst-admission + prefill-determinism gate\n```\n\nThe CUDA engine is validated for **bit-exactness** against reference paths (greedy byte-identical\noutput, `compute-sanitizer` memcheck clean) and benchmarked with the `scripts/` harnesses\n([`pi_bench.py`](scripts/pi_bench.py), [`parity_bench.py`](scripts/parity_bench.py),\n[`benchmark_gem4.py`](scripts/benchmark_gem4.py)) for Gemma, and against llama.cpp/vLLM/`torch`\noracles for Qwen (see the `qwen35-*-test` targets in the [Makefile](Makefile)).\n\n---\n\n## 📌 Status\n\n**Experimental, hardware-specific, no support.** fucina targets exactly one platform — the NVIDIA\nDGX Spark GB10 (Blackwell `sm_121a`) with CUDA 13.0 — and is not portable to other GPUs or\ntoolchains as-is. It is an open research/lab project from **[hikmaai.io](https://hikmaai.io)**,\nprovided **as-is with no warranty and no support commitment**; issues and PRs are handled\nbest-effort. Weights you supply are governed by their own upstream license: Gemma 4 by the\n[Gemma license](https://ai.google.dev/gemma/docs/gemma_4_license); Qwen checkpoints by their\nrespective Qwen license terms on Hugging Face. fucina's own code is Apache-2.0.\n\n- **Code:** [Apache-2.0](LICENSE) · **Third-party notices:** [NOTICE](NOTICE)\n- **Model support:** Gemma 4 12B (dense, GGUF/NVFP4) and Qwen3/Qwen3.5/Qwen3.6 (dense and MoE,\n GGUF/FP8-block safetensors/NVFP4) are both first-class; continuous batching is mandatory for\n Qwen and opt-in for Gemma (see [Continuous batching & paged KV](#-continuous-batching--paged-kv)).\n- **Roadmap:** for Qwen, the current known gaps vs `vLLM` are cold-prompt (turn-1) TTFT and\n aggregate throughput at high concurrency (see [Performance](#-performance) and\n [`docs/qwen35-beat-vllm-plan.md`](docs/qwen35-beat-vllm-plan.md)); `response_format`/`json_schema`\n does not work under continuous batching, so it's unavailable for Qwen today. For Gemma, per-slot\n spec decode inside the batch path is not yet built, and the NVFP4 KV codec / packed 4.5-bit KV\n storage are verified but still opt-in behind a quality gate\n ([`docs/kv-quant-exploration.md`](docs/kv-quant-exploration.md)). Also: harden the experimental\n DiffusionGemma path; an sm_120 (RTX 50-series) port to loosen the single-hardware constraint.\n\n## 🙏 Acknowledgements\n\n- The name **fucina** is Italian for *forge* — a smithy, and figuratively a *crucible of ideas*.\n- **Google** for [Gemma 4](https://ai.google.dev/gemma) and the QAT GGUF / MTP assistant releases.\n- **Alibaba / the Qwen team** for [Qwen3, Qwen3.5, and Qwen3.6](https://huggingface.co/Qwen) and\n their official FP8-block checkpoints.\n- **NVIDIA** for [ModelOpt](https://github.com/NVIDIA/TensorRT-Model-Optimizer) and its NVFP4/FP8\n mixed-precision checkpoint format.\n- **[llama.cpp](https://github.com/ggml-org/llama.cpp)** — the GGUF format, quantized `dp4a`\n kernels, and the `draft-mtp` speculation design that this project measures itself against.\n- **[vLLM](https://github.com/vllm-project/vllm)** — the continuous-batching / paged-KV design this\n project's batching implementation follows, and the head-to-head benchmark target for Qwen.\n- **[unsloth](https://huggingface.co/unsloth)** for the GGUF conversions used here.\n\n\n",
"bytes": 58294,
"sha": "e95f59ecaa80779b2e3f0ae8c433243595d3b4184d5e17caf900f962329cb713",
"repo_slug": "hikmaai-io/fucina",
"fonte": "repo",
"truncated": false,
"api": "https://agentalog.com/api/listings/okf_hikmaai_io_fucina_docs_okf_index_md_7d7a9d15/readme"
}