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

> Research date: 2026-09-15 | Source: github.com/ToddThomson/Mila | Site: mila.toddt.me > Author: Todd Thomson (Port Moody, BC, Canada; independent developer) | License: MIT | Version: v0.20.0-beta.3

This is a structured summary of a long-form Chinese forum investigation into Mila, a C++23 LLM inference library maintained by a single independent developer. The original post is a four-layer verification (model existence, toolchain existence, code authenticity, performance evidence chain). All performance numbers are author-reported and not independently reproduced.

Key stats

| GitHub Stars | Forks | Files (dev) | Spec docs | Commits (since 2021) | Prefill vs llama.cpp | |---|---|---|---|---|---| | 8 | 0 | 975 | 34 | 450+ | 1.14x (still slower) |

One-line verdict (from the reviewer): The project is genuinely real — a non-academic indie developer, five years and one 12 GB GPU, running Qwen3.8-27B in 12 GB VRAM and closing the prefill gap to llama.cpp from 1.95x to 1.14x with hand-written PTX, no CUTLASS. Its weakness is not technical quality but that only one person on Earth verifies these numbers.

Key points

• The gap: A repo with company-release-page polish (badges, roadmap, capability matrix) but 8 stars, 0 forks, and zero external discussion on Reddit/HN/Zhihu/V2EX. The investigation set out to determine whether it is AI-flavored filler or an unnoticed gem — and concludes it is the latter.
• Author background: Todd Thomson previously built TypeScript tooling (tsproject, Ts2Js, TsMinifier) and contributed to Microsoft MSAGL. Mila started as llm (raw C/CUDA training). Explicitly cited inspiration: Karpathy's llm.c and FlashAttention/online-softmax work.
• Types-as-configuration: Device, precision, and quantization policies are template parameters, not config.json. Quantization strategies are empty structs of static constexpr fields. An OperationTraits primary template is *declared but never defined*; a concept OperationSupported via requires { sizeof(...) } makes unsupported combinations a hard compile error. ~48 CUDA specializations exist; adding a quantization format is purely additive.
• PrecisionPlan: Per-role bit widths inside one model — e.g. Qwen: QkvProjection = PerGroupFp4<128> (4.125 b), FeedForwardGateUp = PerGroupCodebook2<32> (2.5 b), DeltaNetGating = NoWeightQuant (BF16, never quantized). 25.62 B params compressed to avg 2.90 bits / 8.65 GiB (the README's "2.82 bits" and "15.1 GiB FP4" figures are flagged as inconsistent documentation).

How 27B fits in 12 GB

• Gated DeltaNet: 48 of Qwen3.8's 64 layers are DeltaNet recurrent layers with a fixed-size state (0.141 GiB FP32) instead of a KV cache that would cost ~6.0 GiB at 32K — a ~42x saving. Trade-off: lossy, non-rollbackable state; the model rejects prompt-prefix reuse.
• Weight tying: Gemma 4 12B shares embedding/lm_head (262144 × 3840), saving ~2.0 GB (docs honestly revise the FP4-case saving to 0.52 GB/table and admit a ~2.0 GB double-count in getMemoryStats).
• Bounded sliding-window ring KV: 40 sliding-window layers keep only 1024 tokens in a ~0.34 GB ring buffer; persistent KV slope drops from 336 KB/token to 16 KB/token (21x flatter).
• Pre-allocation exactness: getRequiredMemory() before build() predicts VRAM byte-for-byte for FP4 and codebook configs; residuals attributed to driver 2 MiB rounding, scratch, and cuBLASLt.

Performance: 1.95x to 1.14x in five documented steps

| Stage | Technique | Throughput @48K (tok/s) | vs llama.cpp | |---|---|---|---| | Start | — | ~1,056 | 1.95x | | +100 | Distributed softmax + ldmatrix | 1,205 | 1.71x | | +101 | Bounded-ring flash on 40 SW layers | 1,372 | 1.50x | | +103 | W4A8-FP8 prefill GEMM (FP8 tensor cores) | 1,763 | 1.17x | | +104 | Split-row FA-2 kernel | 1,817 | 1.14x | | llama.cpp reference | same HW/model/context | 2,063 | 1.00x |

Decode (Flash-Decoding for MQA): 38.65 → 49.09 tok/s vs llama.cpp's 50.3–50.7 — hence README's "within 1.03x". The author consistently reports being slower than llama.cpp, which lends credibility. A seeming contradiction (a spec doc saying "~4x slower than llama.cpp") resolves to a mid-iteration snapshot of a different code path; both kernels shipped.

Two textbook-failure moments

• Head dim 512 won't fit registers: FA-2's row-split explodes; the author instead splits the head dimension across warps, which forces dropping wmma for raw PTX mma.sync.aligned.m16n8k16.f32.bf16.bf16.f32 to get per-row online-softmax rescaling.
• The bug that passed all tests: ldmatrix.x4 PTX missing the .shared qualifier dereferenced shared-memory offsets as generic addresses — parity oracles stayed green because small offsets aliased correct data. New rule: every PTX change must pass compute-sanitizer memcheck.
• Related lessons: an uninitialized scale factor sB produced green per-layer FP8 oracles but garbage generations; a textbook shared-memory K/V staging (16x DRAM traffic cut) bought only 1.2x because the kernel was instruction-issue-bound, not bandwidth-bound — "the profiler falsifies your mental model."

Positioning and governance

• "Mila is raylib, not Unity": a library, not a framework — the app owns main(); no runtime-services middle layer. Three adaptors (Chat, MIS server, future Agentic) split on who closes the loop; MIS is stateless (O(context) per turn) while Chat/Agentic keep warm KV (O(new tokens)).
• Honest cracks: 28 BACKLOG items with 0 DONE; the Docker image has never had a publish build; Python bindings fail to import in-container; a fake-green ldd check; PyPI claims Linux wheels but only win_amd64 exists; test coverage partially restored (24 of ~70 files live); training scope reduced to FP32 GPT-2/MLP (CudaGqaOp::backward throws by design).
• Hardware ceiling: everything validated on one 12 GB RTX 4070; the author requests hardware (RTX 5090) rather than cash. A commit shows an early RTX 5060 Ti (sm_120) native NVFP4 test at 329–361 TFLOP/s.
• Toolchain claims: all 7 version requirements verified real, systematically 1–4 minor versions behind (CUDA 13.0/13.3, VS 2026 ≥18.6.2, CMake ≥4.0, Ubuntu 24.04/26.04, clang-19/GCC 16).
• CLAUDE.md: a 21,850-byte behavioral contract with an AI coding agent — AI never runs git commands, no Co-Authored-By trailers, human-only commit rights, strict edit boundaries. This explains the star/file gap: AI amplified output, but only one human reviews.

Should you use it?

• Yes: Windows + RTX 30/40 series users via the Docker path; ~10 minutes to a Gemma 4 12B FP4 chat at 1.03–1.14x of llama.cpp. A playable, readable, hackable beta.
• No: production use, Linux wheels, pre-Ada GPUs, multi-GPU. The release pipeline has never completed a publish build, and all validation happened on one 12 GB card.

Final verdict

All four verification layers pass: the models exist (Gemma 4 12B, Qwen3.8-27B, both Apache 2.0), the toolchain is real, the code is real (975 files, 34 specs, ~48 trait specializations, two hand-written flash kernels), and the performance numbers have a commit-anchored evidence chain — but all self-reported. Mila is a craft project: its real deliverable is a readable path from main() down to PTX instructions, plus honestly recorded failures. In an era when AI can generate ten polished "high-performance inference engine" READMEs a day, the credibility test is whether a project dares to document how slow it is, where it went wrong, and what it could not verify. Mila passes on all three counts.

---

*All performance figures are author-reported and lack third-party reproduction. Internal documentation inconsistencies (2.82 / 2.90 / 15.1 GiB) are flagged above without reconciliation. Research method: direct source/spec reading + official blog cross-referencing + external toolchain release-page verification.*