8 Stars, 975 Files: One Developer Rewrites LLM Inference from the Metal Up in C++23 (Mila Deep Dive)

This post is a detailed investigation of Mila, an LLM inference runtime built from scratch by independent developer Todd Thomson (British Columbia, Canada): github.com/ToddThomson/Mila | mila.toddt.me | MIT license | v0.20.0-beta.3.

Key points

• The gap: 8 GitHub stars, 0 forks, and zero external discussion (Reddit, HN, V2EX all searched) — yet 975 files, 34 architecture spec documents, 450+ commits since 2021, and ~88 MB of code. The investigation asks: AI-generated bloat, or genuinely unseen craft?
• Four-layer verification passes: the models (Gemma 4 12B, Qwen3.8-27B, both Apache 2.0), toolchain versions (7/7 real, though lagging 1–4 minor releases), the code (real CUDA kernels, ~48 OperationTraits specializations), and the performance claims (fully documented but self-reported).

Technical highlights

• Types-as-configuration: device, precision, and quantization strategy live entirely in template parameters. Mixing CPU and CUDA tensors is a compile error, not a runtime surprise. Unsupported operation/policy combinations fail via an undefined primary template (OperationTraits), making illegal configurations physically unexpressible.
• Per-role precision (QwenPrecisionPlan): FFN gate-up at 2.5 bits, gating layers and embeddings at BF16 — 25.62 B parameters compressed to an average 2.90 bits / 8.65 GiB (README's 2.82 and 15.1 GiB figures describe different configurations; flagged as a doc inconsistency).
• 12 GB fits 27B via: 48 of 64 layers being Gated DeltaNet recurrent layers (fixed-size state: 0.141 GiB vs ~6 GiB for hypothetical full attention at 32K, a ~42× saving), weight tying (~2 GB saved on Gemma 4 12B), a bounded sliding-window ring KV buffer (flat ~0.34 GB regardless of context), and exact memory prediction before allocation (getRequiredMemory() matched measured values byte-for-byte).

Performance evidence chain

Prefill on RTX 4070, Gemma 4 12B @48K, tracked commit-by-commit from 1.95× slower than llama.cpp to 1.14× slower (~1,056 → 1,817 tok/s vs llama.cpp's 2,063), via distributed softmax, bounded-ring flash on sliding-window layers, an FP8 tensor-core W4A8 GEMM, and split-row FA-2. Decode reached 49.09 tok/s vs llama.cpp's 50.3–50.7 (the README's "within 1.03x"). Notably, the author consistently reports *still being slower* — a credibility signal.

Instructive engineering war stories

• Head dim 512 overflows registers: textbook FA-2 row-splitting fails; the fix is head-dimension splitting across warps, forced down to raw mma.sync.aligned.m16n8k16.f32.bf16.bf16.f32 PTX for per-row online-softmax rescaling.
• A bug that passed all tests: a missing .shared qualifier on ldmatrix PTX aliased to correct data at small smem offsets — green only via undefined behavior, caught by compute-sanitizer. New rule: every PTX change must pass memcheck, parity alone is insufficient.
• An uninitialized scale factor passed per-layer FP8 oracles but produced garbage across 48 layers — caught by initcheck.
• The most counterintuitive "failure": textbook shared-memory K/V tiling cut DRAM traffic 16× but bought only 1.2× — the kernel was instruction-issue-bound, not bandwidth-bound (L2 hit rate 99.7%). Lesson: "the profiler falsifies your mental model."

Honest cracks

• Release pipeline never succeeded: 28 backlog items, 0 DONE; Docker image never published; Python binding unimportable in-container.
• Fake-green gates: an ldd check that passes when it finds nothing; PyPI claims Linux wheels but only win_amd64 shipped.
• Test coverage partially restored (24 of ~70 test files live); training limited to FP32 GPT-2/MLP; GQA backward throws by design.
• Hardware ceiling: everything validated on one 12 GB RTX 4070; FP4 weights for 27B+ models don't fit. The author requests hardware, not money — specifically an RTX 5090 (sm_120) to test Blackwell's native FP4.

The AI-workflow reveal

A 21,850-byte CLAUDE.md acts as a behavioral contract for AI coding assistants: humans-only git commits, no abbreviations in identifiers, spec docs as "design of record," and "Done means deleted." This explains the star/file-count gap: AI amplified output, but no one shares the review burden — so all discipline is encoded as rules for the AI.

Verdict

The project is real — arguably *too* real for its own visibility. It's a craft project (the author's own framing: "Mila is raylib, not Unity"), explicitly declining throughput leaderboards. Its deliverable is a readable path from main() to PTX instructions, plus honestly documented errors. Suitable for Windows + RTX 30/40-series users via Docker (Gemma 4 12B FP4 chat in ~10 minutes); not recommended for production, Linux wheels, or multi-GPU. As the author concludes: in an era when AI can generate ten polished inference-engine READMEs a day, the test of authenticity is whether a project dares to write down how slow it is, where it's wrong, and what it couldn't verify.

Method note: All performance figures are author self-reported without third-party reproduction. Internally inconsistent figures (2.82 / 2.90 bits / 15.1 GiB) are flagged, not reconciled. Docker Hub data could not be fetched and comes from the project's own site.

Original source: repo dev branch, 34 spec documents, 11 blog posts at mila.toddt.me, PyPI (mila-llm 0.20.0b3), and HuggingFace.