Inference engine · NVIDIA DGX Spark
Ferrite is an LLM inference engine for NVIDIA DGX Spark (GB10), written in Rust from the HTTP server down to the GPU kernels. It serves Qwen3.8-27B on one Spark or across two connected over RoCE, behind an OpenAI-compatible API with streaming, thinking, tool calls, a prompt cache and 128K context.
Status: experimental. One model (Qwen3.8-27B), one GPU (GB10, sm_121a).
At most 4 requests run at once with DFlash2 drafting; more wait in a queue. It works, but expect rough edges.
Issues and measurements from other Sparks are welcome.
Why it exists
Ferrite started as an experiment by a TensorFold contributor: how far down the stack can Rust go in a real inference engine, all the way into the GPU kernels, and what does it cost? Much of what makes Ferrite work was learned from TensorFold and exllamav3 first.
If you want a mature engine that serves many models, use TensorFold. Ferrite is a focused, single-model look at the Rust question.
Rust: every GPU kernel, the CUDA driver bindings, the model and scheduler, speculative decoding, the two-node fabric (RDMA over libibverbs), the tokenizer and the HTTP server.
Not Rust: NVIDIA's ptxas assembler (at build time), the CUDA driver itself,
libibverbs, and NCCL for two-node prompt batches. The full ledger, with reasons and replacement plans, is in
NON-RUST.md.
Design
The kernels are ordinary Rust crates. rustc's nvptx64-nvidia-cuda backend turns them into
PTX, NVIDIA's ptxas assembles that for sm_121a, and the cubins are embedded in the server
binary and launched through a thin Rust binding to the CUDA driver.
Decode reads Qwen3.8-27B as an EXL3 quant at 3.0 bits per weight (a mixed-width quant with a 5-bit head) at about 90% of GB10's memory bandwidth.
The matrix-vector kernels batch up to 48 rows into one weight decode, and every row stays bit-identical to single-row decoding.
Prompts run through NVIDIA's NVFP4/FP8 checkpoint on tensor-core GEMMs written in Rust (block-scaled
mma.sync: NVFP4 with ue4m3 block scales, FP8 e4m3), with the activation quantizers on the
GPU too.
Without that checkpoint, Ferrite rebuilds bf16 prompt weights from the EXL3 quant.
Qwen3.8's linear-attention layers carry a recurrent state per request. Prompt rows go through them in chunks of 32 rows, using the chunkwise WY / UT form of the delta rule on TF32 tensor cores (a 3xTF32 split for fp32 accuracy), and the state is written back for decoding.
Lossless speculative decoding. The DFlash2 block drafter proposes tokens that are verified as best-first trees in one pass of the big model.
Copy rows propose the continuation of earlier text, which pays off most on code edits. The model's own MTP head is the fallback.
With two nodes, each Spark holds half of every layer. Small reductions go over RDMA through a Rust proxy kernel, and NCCL takes the large prompt batches. Node A runs the server, node B follows.
Chat completions with streaming, thinking on or off, tool calls, 128K context and up to 4 requests at once (more wait in a queue). Any OpenAI-compatible client works: Open WebUI, Continue, Cline and other coding agents.
A request's output does not depend on what else is running, and speculation never changes the text. Both are checked before every merge.
Each conversation is snapshotted at turn boundaries and resumes from the longest match. A conversation that loses its slot is parked in a memory pool, and a system prompt shared by several agents is computed once.
Prompts are cut at the same points with or without a resume, so a resumed reply is the same, token for token, as one computed in full.
Measured
Two DGX Sparks connected over RoCE, Qwen3.8-27B, greedy decoding, fixed prompts. Medians of 3-5 runs. Generation figures are tokens per second while generating; prompt processing is listed separately. End-to-end figures that include the time to the first token are in docs/benchmarks.md.
| Workload | Ferrite |
|---|---|
| Code, 4 requests (aggregate) | 473 |
| Code, 1 request | 273 |
| Code, 1 request, 32K context | 238 |
| Code edits, 1 request (rewrite a file it was shown) | 582 |
| Code edits, 4 requests (aggregate) | 481 |
| Prose, 1 request | 76 |
| Prompt processing, 2K / 32K context | ~3,550 / ~2,800 |
| KL divergence vs BF16 (code / text / 16K context), lower is better | 0.021 / 0.040 / 0.024 |
Speed by context length. Medians, with the minimum and maximum as bars.
One request's generation is timed from its first token to its last (500 tokens); four requests give their combined rate after the last one's first token. Prompt processing is timed to the first token with a cold prompt cache, and accuracy is the KL divergence of the next-token distribution against the BF16 model over fixed code, text and long-context sequences. The method, the clients and how to reproduce them are in docs/benchmarks.md.
Agents
Four coding agents at once on two Sparks, sharing a ~13K-token system prompt and running 2-3 turns
each (tools/bench/agents.py), with the prompt cache on instead of off. The cache is on by default;
FERRITE_CACHE=0 turns it off. Its park pool is 12 GB on one Spark and 24 GB on each of two
(FERRITE_PARK_GB overrides), and it shrinks if needed to keep 20 GB of memory available.
How it works, what it guarantees and its limits are in docs/prompt-cache.md.
Quickstart
Needs a DGX Spark (GB10) with DGX OS, NVIDIA driver 580+, at least 95 GB of available memory and
38 GB of disk. One command takes a fresh Spark to a running server: it checks the machine, installs Ferrite and
the models, starts the server and tests it. No sudo, and it can be run again.
ferritecurl -fsSL https://raw.githubusercontent.com/BHCC2025/ferrite/main/scripts/install.sh | bashInstall with your AI agent. Paste this into your coding agent on the Spark:
Install Ferrite on this DGX Spark by following https://github.com/BHCC2025/ferrite/blob/main/AGENTS.md exactly; stop and ask me if any check fails.
Or step by step, from the release tarball (no Rust or CUDA toolchain needed):
# 1. unpack the release and check the machine
tar xzf ferrite-<version>-dgx-spark-aarch64.tar.gz
cd ferrite-<version>-dgx-spark-aarch64
scripts/check-deps.sh
# 2. download the models (about 38 GB) into $FERRITE_MODELS
export FERRITE_MODELS=$HOME/models
scripts/fetch-models.sh
# 3. serve an OpenAI-compatible API on :8000 (model id "spark")
scripts/serve.sh
curl -s localhost:8000/v1/modelsSame Ferrite folder and model folders on both nodes, passwordless ssh from node A to node B, NCCL on both. Then on node A:
FERRITE_PEER=<node-b> tools/serve-tp2.shThe two-node scripts detect the RoCE devices, interface and GID the nodes share;
tools/site.sh prints every resolved setting.
Stable Rust for the host crates, the kernels' pinned nightly (nightly-2026-11-04)
with rust-src, llvm-tools and llvm-bitcode-linker, and CUDA 13's
ptxas. The check script installs the Rust toolchains.
scripts/check-deps.sh --build --install
cargo build --release -p fe-server --features ferrite -p fe-cliPoint any OpenAI-compatible client at http://<spark>:8000/v1. The full guide, with every setting,
client examples and troubleshooting, is in
docs/quickstart.md.
| Role | Hugging Face repository |
|---|---|
| Decode weights: EXL3 mixed 3.0 bpw (5-bit head) | BHCC2025/Qwen3.8-27B-EXL3-3.0bpw-mix |
| Prompt weights: NVFP4/FP8 (optional) | nvidia/Qwen3.8-27B-NVFP4 |
| Draft model: DFlash2 | z-lab/Qwen3.8-27B-DFlash2 |
All three are Apache-2.0 and derive from Qwen/Qwen3.8-27B (Apache-2.0). Without the NVIDIA checkpoint, Ferrite processes prompts with bf16 weights rebuilt from the EXL3 quant, at about half the prompt speed.
Credits
Ferrite stands on a lot of other people's work. CREDITS.md lists every project and paper it drew on and exactly what it took from each. The largest debts: