8 颗星、975 个文件：一个人用 C++23 从金属层重写 LLM 推理（Mila 深度研究）

研究日期：2026-09-15 ｜ 来源：github.com/ToddThomson/Mila ｜ 官网：mila.toddt.me
作者：Todd Thomson（加拿大 BC 省，独立开发者，一人维护）｜ 许可：MIT ｜ 版本：v0.20.0-beta.3（树内已推进至 0.20.0-rc.1+10）
取材：仓库 dev 分支源码与 34 篇规格文档、官网与 11 篇博客、PyPI / HuggingFace / Release API，以及两组外部交叉核验。所有数字逐条指认出处，未能证实的已标注。

一句话结论：这是一个真货，而且真得有点过分。一个没有任何学术背景的独立开发者，用五年时间、一张 12 GB 显卡，把 Qwen3.8-27B 塞进去跑起来，并把「prefill 落后 llama.cpp 1.95 倍」一路徒手追到 1.14 倍——全程手写 PTX，不碰 CUTLASS。它的问题不在技术成色，而在于：全世界只有他一个人在验证这些数字。

一、第一个钩子：8 颗星与 975 个文件的巨大裂口

先说最刺眼的反差。这个仓库的 README 写得像一家公司的产品发布页：徽章、路线图、能力矩阵、性能战报、四语言上手路径，一应俱全。按常理，这种完成度至少该有几百颗星。

实际数据：

八颗星、零分叉、中文与英文技术社区零讨论——所有外部链接最终都绕回作者自己的网站。这不是「火了但没人 fork」，这是根本没有被看见过。

那么问题来了：这到底是「文档写得漂亮、代码注水」的 AI 味产物，还是「真材实料但无人知晓」？我把它拆成四层验证：模型是否真存在、工具链是否真存在、代码是否真写了、性能数字是否有证据链。

二、来路考：一个 TypeScript 工具链老兵的中年转型

作者 Todd Thomson，住加拿大不列颠哥伦比亚省 Port Moody，公司栏填「Achilles Software」，个人简介四个词：Kind of a Nerd。他的个人站自述很实在：

「我写开源开发者工具——编译器、构建工具链，现在是机器学习运行时。」以及一句关键的：「我是一个独立开发者 (a single developer)。」

转型轨迹很有意思：做 CUDA 之前，他写了好些年 TypeScript 工具链——tsproject、Ts2Js、TsMinifier，还有 jQuery Mobile 插件，并且给微软的 MSAGL（自动绘图布局库）贡献过 UWP 版本。这些项目如今都已归档休眠。

早期 Mila 的仓库名就叫 llm，是「raw C/CUDA 训练」。也就是说：他不是从 PyTorch 转过来的，他是做编译器和构建工具出身的人，从金属层往上爬。这解释了这个项目最鲜明的气质——对「显式」近乎偏执的追求。

值得注意的是：搜索「Todd Thomson」会大量命中一位华尔街金融人士（花旗前 CFO），与本仓库无关，属同名干扰。没有找到他的 LinkedIn、X/Twitter，也没有找到任何学术机构归属或论文。

致敬清单里的一个彩蛋

ATTRIBUTIONS.md 只认一个灵感源：Karpathy 的 llm.c（「通过从零构建来理解它」）。技术债主要给 FlashAttention/FA2/Flash-Decoding（Tri Dao 等）和 Milakov & Gimelshein 的 online softmax。llama.cpp 只作为性能标尺出现，不是代码来源。

最后一条致谢让人会心一笑——一本童书，《迈克·穆利根和他的蒸汽铲》。「no hidden engine」（没有隐藏引擎）这个自我定位，出处在这里。

三、信仰：类型即配置，把「运行时惊喜」变成「编译期错误」

Mila 的所有设计，都从一句话长出来：混用 CPU 张量和 CUDA 张量，必须是一个编译错误，而不是运行时的意外。

这句话听着朴素，做起来是另一回事。市面上的框架普遍做法是：你写个 config.json，运行时读进来，然后一路 if (device == "cuda") 分支派发。Mila 不这么干——它把设备、精度、量化策略全部塞进模板参数：

// 没有 config 对象，没有运行时派发。类型本身就是配置。
using QuantizedProjection =
    Linear<DeviceType::Cuda, TensorDataType::BF16, PerGroupFp4<128>>;

打个比方：这好比电源插头。三脚插头插不进两脚孔——不是因为它会提醒你「警告：电压不匹配」，而是它物理上就插不进去。Mila 要把所有「本该在运行时崩」的错误，提前变成「编译期就过不去」。

策略是空壳 struct，全部编译期常量

量化策略不是对象，是只带 static constexpr 字段的空结构体：

export template<int kGroupSize = 128> struct PerGroupFp4 {
    static constexpr bool kIsQuantized = true;
    static constexpr TensorDataType kStorageDtype = TensorDataType::UINT8;
    static constexpr bool kPerChannel = false;
    static constexpr int kQuantizationGroupSize = kGroupSize;
    static constexpr bool kIsFp4E2M1 = true;
    static constexpr int kStorageBitsPerElement = 4;
};

OperationTraits：一张「没实现就编译不过」的分派表

核心机关在这里。主模板只声明、不定义：

export template<OperationType TOp, DeviceType TDeviceType,
                TensorDataType TPrecision, typename TPolicy = void>
struct OperationTraits;   // 注意：无定义

export template<OperationType TOp, DeviceType TDeviceType,
                TensorDataType TPrecision, typename TPolicy = void>
concept OperationSupported =
    requires { sizeof( OperationTraits<TOp, TDeviceType, TPrecision, TPolicy> ); };

CUDA 后端里约 48 条特化，形如：

template<> struct OperationTraits<OperationType::LinearOp, DeviceType::Cuda,
                                 TensorDataType::BF16, PerGroupCodebook2<32>>
{ using type = CudaLinearOp<TensorDataType::BF16, PerGroupCodebook2<32>>; };

妙处在于：**没有 kernel 支撑的组合，::type 就落到那个「只有声明没有定义」的主模板上，于是报不完整类型错误。**CPU-only 构建下，CUDA 那批特化压根不存在（用 #ifdef MILA_HAS_CUDA 隔离），所以「CPU 上跑 FP4」这种组合在物理上无法被表达。新增一种量化格式是纯增量——加一个特化块，不需要注册器。

PrecisionPlan：同一个模型里，不同角色不同位宽

这是我认为全项目最优雅的一处。传统量化是「整个模型统一 4bit」。Mila 说：凭什么？FFN 的 gate_up 能扛 2bit，门控投影却一点都不能动。于是它做了一个按角色命名的 typedef 集合：

export struct QwenPrecisionPlan {
    using QkvProjection              = PerGroupFp4<128>;          // 4.125 b
    using FeedForwardGateUp          = PerGroupCodebook2<32>;     // 2.5   b
    using FeedForwardDown            = PerGroupCodebook3<64>;     // 3.25  b
    using DeltaNetQueryKey           = PerGroupCodebook3<64>;
    using DeltaNetValueGateOutput    = PerGroupCodebook2<32>;
    using DeltaNetGating             = NoWeightQuant;   // BF16，永不量化
    using EmbeddingTable             = NoWeightQuant;
    using LanguageModelHead          = PerGroupFp4<128>;
};
static_assert( DecoderPrecisionPlan<QwenPrecisionPlan> );
static_assert( DeltaNetPrecisionRoles<QwenPrecisionPlan> );

位宽算术也干净：2 + 16/32 = 2.5、3 + 16/64 = 3.25、4 + 16/128 = 4.125——即「码本位数 + 每组 16 bit 的缩放因子摊薄」。最后 25.62 B 常驻参数压到 平均 2.90 bit、8.65 GiB。

一处诚实的瑕疵：README 写「平均 2.82 bits」，源码 Qwen.PrecisionPlan.ixx 与 Qwen3.8.md §5 都写 2.90 avg / 8.65 GiB，README 另一处又写 FP4 版 15.1 GiB。三个数字并存。我的判断：2.82 是质量门实验表头里某一档的标签，2.90 是名义平均，15.1 GiB 是 FP4 版——是两组不同的配置，但文档没有把话说清，读者容易混淆。

四、显存术：12 GB 卡是怎么塞下 27B 的

作者的验证机器只有一张 RTX 4070（12,282 MiB）。在这个约束下跑 Qwen3.8-27B，靠的是四件事叠加。

① Gated DeltaNet：48 层不写 KV cache

这是省得最狠的一招。Qwen3.8-27B 是混合架构：64 层里 48 层是 Gated DeltaNet 循环层，只有 16 层是全注意力（每 4 层来一次）。

为什么这很关键？传统 Transformer 每生成一个 token，都要把 Key/Value 追加进 KV cache，序列越长越膨胀，是显存杀手。而 DeltaNet 层维护的是固定大小的循环状态——不管你说了一千字还是十万字，它那份状态就这么大。

也就是说，在这一半层数上，本来要花 6.0 GiB 买的东西，现在 0.141 GiB 就解决了，省下约 42 倍。

代价作者写得很清楚：循环状态是对历史的有损摘要，无法回滚。所以 Qwen「彻底拒绝 prompt 前缀复用」——任何想复用前缀的服务器，都得把这条当作模型属性来适配，而不是当成一个可以重试的失败。

② 权重绑定（weight-tying）：省掉一整张词表

Gemma 4 12B 的 embedding 与 lm_head 共享同一张 262144 × 3840 的表，直接省下 约 2.0 GB。

不过文档中有一处很实在的自我修正：FP4 构建里这张表实际走 PerChannelFp8<>，量化表 0.53 GB 对 BF16 的 1.05 GB，每张只省 0.52 GB，而不是最初记录的 0.79 GB；而且 getMemoryStats 因 config 与 metadata 双源不一致，在 Gemma 4 12B 上重复计数了约 2.0 GB。这种把自家账目改错的过程原样写进文档的做法，比很多项目的 README 可信得多。

③ 有界滑窗环形 KV

Gemma 4 是 40 个滑窗层 + 8 个全局层。滑窗层只需保留最近 1024 个 token，用一个环形缓冲：

capacity = min(T, window + prefill_chunk − 1) = 1087 slots   // 恒定 ~0.34 GB

持久 KV 斜率从 336 KB/token 降到 16 KB/token，平缓了 21 倍。

④ 分配之前就能算出要多少显存

这一条最能体现工程成熟度：构造计算图不占用任何设备字节，在 build() 之前调用 getRequiredMemory()，每个组件按真实配置自报。FP4 的每参数字节数是 0.5 packed + 4/128 = 0.53125。

结果是：Gemma 4 12B FP4 @8192、Llama 8B FP4、Qwen 27B FP4、codebook 2-3bit，预测值与实测值逐字节相等。残差被归因到驱动 2 MiB 舍入（0.02–0.73 GiB）、scratch（0.07–0.25 GiB）、cuBLASLt 4 MiB。

文档里有一句我很喜欢的话：显存预算里约 1176 MiB 的基线不是 CUDA context，而是桌面合成器。作者拒绝建模它，理由是「预测它等于预测用户的浏览器」。

五、性能战报：1.95× 到 1.14× 的十轮会战

这是全篇最硬的部分，也是我一开始最怀疑的部分。README 宣称「prefill 达到 llama.cpp 的 1.14 倍」「decode 在 1.03 倍以内」——这种话谁都会写。

但我找到了完整证据链。作者把整个过程写成了博客 Writing a FlashAttention Prefill Kernel for Gemma 4。文中明确写着起点：

「十轮之前，Mila 的 Gemma 4 12B prefill 在 RTX 4070 上、48K 上下文，比 llama.cpp 慢约 1.95 倍。」

终点是 0.20.0-alpha.6+104 版本的 1.14×。中间五步，每一步都有 commit 号和吞吐数字：

这组数字的可信度为什么高：每一步都锚定 commit 号（+100/+101/+103/+104），参照组是同机同模型同上下文的 llama.cpp，而且作者承认自己仍然更慢。1.14× 是「还在落后 14%」的意思，不是「快了 14%」。一个要吹牛的人不会这么写。

decode 侧：同样是自报，但同样留有对照

另一篇 Flash-Decoding for MQA in CUDA 给出：Gemma 4 12B FP4 @32K 上下文，从 38.65 tok/s 提到 49.09 tok/s；作者自己实测 llama.cpp 同期为 50.3–50.7 tok/s。49.09 / 50.7 ≈ 1.03——这就是 README 那句「within 1.03x」的来历。依然是承认落后。

过程中的分解很漂亮：nsys 把单 token 的 24.9 ms 拆成 FP4 matvec 15.17 ms、GQA attention 4.63 ms、RMSNorm 2.47 ms、lm_head 2.16 ms。然后逐个击破——用 PRMT 字节置换替换常量内存查表，把 DRAM 峰值占用从 63–78% 拉到 79–90%；GQA decode bucket 从 4.63 ms 压到 0.37 ms（12.5 倍）。

我原本以为抓到了一处「吹牛」，结果被反证

调研中我曾发现一条矛盾，值得记录：规格文档 GqaFlashAttention.md 里白纸黑字写着 "Measured ~4x slower than llama.cpp"，且 use_flash_prefill_ 仍为 false；而 README 写的是 1.14×。看起来像是 README 拔高了数字。
核对博客原文后，结论反转：两边的数字都对，指的是不同的东西。规格文档是开发日志，记录的是 kernel 迭代中期的快照（cuBLASLt 8K prefill 6465 ms vs Flash Iter2 19068 ms，确实慢约 3 倍）；README 的 1.14× 是最终启用的路径。而且最终两个 kernel 都交付了：head dim 512 用 HS-split 给 8 个全局层，head dim 256 用行切分 FA-2 给 40 个滑窗层——所以 README 说「滑窗层」、规格说「全局层」，两个说法都成立。这条我判作者无过。

六、两个「教科书失效」的时刻

这两段是整份材料里最有 teaching value 的部分，值得单独摘出来。

① head dim 512 装不进寄存器

标准 FlashAttention-2 的做法是：把 query 行切给不同 warp，输出 tile O[Br × HS] 留在寄存器里。教科书默认 head dim ≤ 128。可 Gemma 4 的全局层 head dim 是 512——那个累加器要 256 个 float 每线程，直接溢出，不可挽回。

作者的解法是换个刀法：不切行，切 head 维度。第 w 个 warp 认领输出列 [w·64, (w+1)·64)，QK 变成一次 split-K 收缩加一个极小的共享内存归约，PV 写互不重叠的列切片，warp 间无需同步。

这又逼出第二层：因为 online softmax 要对张量核累加器做逐行 rescale，而 wmma 这套便利 API 把 fragment 到 (row, col) 的线程映射藏了起来，根本表达不出来。只能下沉到裸 PTX：

mma.sync.aligned.m16n8k16.f32.bf16.bf16.f32

mma.sync 的布局是有文档的，所以每个线程知道自己手里是哪两行，能在每次累加前施加正确的 alpha。作者的注解很实在：「这就是 CUTLASS/FA-2 的做法，手写的。」

② 那个通过了所有测试的 bug

加了 ldmatrix.x4 之后，parity oracle 通过、聊天输出连贯——然后 ProfileModel 在大规模下崩了，报 cudaErrorIllegalAddress。

compute-sanitizer 揪出真凶：ldmatrix 的 PTX 少了 .shared 限定符，于是共享内存偏移被当成通用地址去解引用。在小的 smem 偏移上，这个错误地址恰好别名到了正确的数据，所以 oracle 一直绿——绿在未定义行为之上。

正确写法是 cvta.to.shared.u64 加指令上的 .shared。作者由此立了新规矩：每一处 PTX 级改动，都必须过 compute-sanitizer memcheck，光 parity 通过不算数。

还有一条同源教训：验证生成，而不是验证 oracle

FP8 prefill 的初版在 5e-2 容差下逐层 oracle 全绿，但跑满 48 层后生成的是垃圾。原因是一个未初始化的缩放因子 sB 在量化器填好之前就被读了——单层误差看不出来，跨层累积就爆了。数值 oracle 看不见它，compute-sanitizer --tool initcheck 一眼看穿。

最反直觉的一次「失败」

作者花了大力气做共享内存 K/V 切分——教科书里真正的 flash 算法——把全局 K/V 流量砍了 16 倍。结果只买来 1.2 倍。

ncu 告诉他为什么：这个 kernel 从来就不是 DRAM 瓶颈（L2 命中率 99.7%）。标量「每 warp 一行」的写法每做一次 FMA 就要发一次共享内存 load，1 FLOP 对 2 字节的比例，卡死在指令发射速率上，不是带宽上。切分把 load 从 L2 挪到了 shared，但没有减少 load 的次数。整条标量路线在结构上就是死路，再怎么调占用率、修 spill 都逃不出去。唯一的出路是张量核。

作者总结：「分析器会证伪你的心智模型。」 他们以为自己是内存受限，实际是指令发射受限，然后是延迟受限，然后是占用率受限——每一阶段的瓶颈都不一样，只有 ncu 能说清是哪一种。

七、定位：Mila 是 raylib，不是 Unity

这是我最欣赏的一份规格文档里的一句话。作者在 MilaProductFamily.md 里反复强调：

「Mila 是一个库，不是框架：应用程序拥有 main()、循环和工具；Mila 让模型成为它们里面的一个普通 C++ 对象。
（那个诱人的类比——『推理的游戏引擎』——指错了方向：引擎拥有循环，然后回调你的代码。Mila 是 raylib，不是 Unity。）」

由此推出一条边界铁律：runtime 只接受「模型固有的、且对消费者无感知的」东西。不存在一个「runtime services」中间层去管会话、历史、prompt 拼装——因为那正是一个隐藏执行引擎一次一个便利类地长出来的方式。

三个 adaptor，按「谁合上循环」来分

关键的结构性优势写在文档里：MIS 是无状态的，客户端每轮重发完整历史，所以每轮都是 O(context)；而 Chat 与未来的 Agentic 共享一个进程内的 agent core，KV cache 全程保温，每轮只 prefill 增量，是 O(new tokens)。这是「有序列化边界的架构结构性地拿不到」的东西。

作者甚至设计好了怎么量这件事：同一个模型、同一套工具、两个循环对打，报告每轮 prefill token 数——O(new tokens) 与 O(context) 的差异要么显形，要么不显形。

一个已经发生的架构教训

规格文档记录了一次真实的漂移：Gemma 的工具调用语法被写了三遍（Chat 里一份 C++，MIS 里一份 Python，还有一处），副本漂移得足够远——Chat 把工具声明渲染成 JSON 数组，而 MIS 渲染的是模型实际训练时的原生声明形式。同一个模型，两种实质性不同的 prompt，取决于用户手里拿的是哪个 adaptor。

作者的判断是对的：这不是打个补丁的 bug，这是分层在告诉你代码放错了地方。语法属于 runtime，adaptor 只做投影。

八、诚实的裂缝：这个项目现在真实的样子

吹完了，说问题。而且这些问题比通常的「还在 beta」要具体得多——好在大部分是作者自己写在文档里的。

裂缝一：发布管道一次都没跑通

最需要警惕的一条。BACKLOG 显示 28 项、0 个 DONE（规则是「Done 意味着删除」），4 个进行中的条目有 3 个卡在容器发布链：Docker runtime 镜像从未有过一次 publish 构建；镜像内 Python binding 无法 import（No module named 'mila._mila'）；干净树仍不可复现。

裂缝二：假绿的门禁

这类问题最危险，因为它让你以为一切正常：

• ldd 在文件缺失时 grep 不到目标，于是脚本打印 「Shared library check passed」——检查通过是因为它什么都没检查到。
• PyPI 上声明了 Linux wheel，但实际只发过 win_amd64。
• 容器默认 MILA_CUDA_ARCH=native，而发布需要 80;86;89;90;120。

裂缝三：测试覆盖只恢复了一部分

ROADMAP 写得很坦白：第一年是测试驱动的，但推理时代的重构把测试大量注释掉了，只剩约 70 个文件里的 24 个还活着。这次发版是「恢复」，其中 14 个文件从未重新启用，量化加载时白盒与 Llama 推理路径的覆盖根本没写。而且原来的测试套件是只前向的——训练样本跑的每一个 backward() 覆盖率都是零。

裂缝四：训练能力被大幅缩水

v0.20 的训练范围只有 FP32 的 GPT-2 / MLP。文档明说：

• CudaGqaOp::backward 按设计直接抛异常（GQA 反向不存在）
• BF16 训练路径在 0.20.0-beta.2+16 之前编译不过、链接不过，且初始化器会越界写自己的缓冲区
• Llama 微调、GQA 训练、BF16 训练全部推到未来版本
• 连 Linear::backward 在量化分支里也是运行时 throw std::logic_error——不是编译期拒绝

裂缝五：硬件，唯一的硬瓶颈

SPONSORING.md 写得非常直白：全部开发与验证机器只有一张 12 GB RTX 4070。Qwen 3.8-27B / Muse Glimmer 30B 的 FP4 权重就要 13.5 / 16 GB，「在 Mila 支持的任何精度下都加载不进来」。

他只要硬件，不要钱——明写不接受 GitHub Sponsors、不收现金；头号请求是一张 RTX 5090（32 GB, sm_120），因为在 Ada 上 FP4 只是存储格式，Blackwell 的原生 FP4 算路他根本没法测。

（有意思的是，commit rc.1+3 显示他已经在 5060 Ti（SM120）上实测过 cuBLASLt 原生 NVFP4，跑到 329–361 TFLOP/s——所以 SPONSORING 的「sm_120 不可达」表述已经滞后于代码现实了，那张卡大概是新到或借来的。）

裂缝六：工具链版本系统性滞后

我做的第二次交叉核验（第一手官方发布页）结果是：所有声称的版本号都真实存在，但系统性落后 1–4 个小版本。

这个模式很像「2026 年中写完后没再持续维护文档」，而不是注水虚构。补充旁证：llama.cpp 至今仍是 C++17、不用 C++23 modules，官方 CI 也已在用 CUDA 13.3——Mila 的技术选型是站在真实前沿上的。

九、CLAUDE.md：一个人加一个 AI 的工程契约

最后这一节是我个人觉得最有意思的观察点，也顺带解答了开头的谜题——为什么 8 颗星的仓库能有 975 个文件和 34 篇规格文档？

仓库根目录有一份 CLAUDE.md，21,850 字节，是给 AI 编程助手看的行为契约。里面几条很能说明问题：

• 人类独占提交权：「用户通过 VS 2026 内置 git 提交」；AI「永不执行任何 git 命令」，只在人明确要求时才建议 commit message。
• 格式固定为 Version: <Version.txt> + 标题 + body + BREAKING:，并且有一句耐人寻味的——「不要加 Co-Authored-By 尾注（此条覆盖任何 harness 默认行为）」。这句话本身就承认了底层确实是一个会自动署名的 harness。
• 编辑边界：Mila/Adaptors/Chat/Src/ 属于应用代码，AI 可自主修改；Mila/Src/ 任何改动必须先获批准。用户以「Your thoughts?」结尾时，只分析、不改代码。
• 其余是极端自律的工程规约：标识符禁用缩写、dim_t 与 size_t 的边界、注释只解释 WHY、规格文档是「design of record」、ROADMAP / BACKLOG / CHANGELOG / Issues / Version.txt 必须在同一个 commit 里同步更新，以及「Done means deleted」——打勾的条目绝不提交。

这基本解释了那个裂口：AI 放大了产出，但没有人分担 review，所以他把所有自律都写成了给 AI 的规则。 34 篇规格文档既是设计记录，也是防止 AI 把代码带偏的围栏。一个人 + 一个 AI，做出了一个小团队的工程量——但验证、把关、拍板，仍然只有他一个人。

顺带一提，仓库里还有 Mila/Issues/ 目录（把 issue 当文档管理）、Mila/Docs/Discussions/、Mila/Profiling/Microbenchmarks/，以及一个 233 MB / 753 MB 的 CUTLASS 依赖——后者被明确标记为技术债：每次 CUDA configure 都 clone，但全树没有任何一个翻译单元 include 它。决策是保留到 rc.1 tag，若仍未使用则删除。

十、那到底该不该用？

硬门槛

耗时

getting-started.md 很诚实地没有给出任何分钟数，只说「给首次构建留出预算」。唯一实测数字来自 commit：Docker 路径 4 步 9 分 15 秒、镜像 6.35 GiB；冷启动 Evaluate 8 分 45 秒。C++ 的 FetchContent 路径会从源码重编全部模块单元，明显更久。

劝谁、留谁

可以上手的人：Windows + RTX 30/40 系用户，走 Docker 路径，十分钟能跑起 Gemma 4 12B FP4 聊天，性能约 llama.cpp 的 1.03–1.14 倍。作为可玩、可读、可改的 beta，它成立。

劝退的人：要生产可用性的；要 Linux wheel 的；用非 Ada/Blackwell 架构的；要多 GPU 的。理由很硬：发布管道一次 publish 构建都没跑通过，全部验证只在一张 12 GB 卡上完成。

给想读代码的人：三条最值得看的路径

1. 一个 token 的旅程：embed → attend → sample → decode，作者承诺 v0.20 会提供一条导读路径，让人能独立跟着源码走完。
2. 那两篇 kernel 博客：FlashAttention Prefill 与 Flash-Decoding for MQA。它们把「为什么教科书在 head dim 512 上失效」「为什么 oracle 全绿还能崩」讲得比论文还清楚。
3. 规格文档里的自我限定：尤其是 Qwen3.8.md 末尾那几句——「它是对 FP4 的比值，不是对 BF16 的；不要把这个比值当成量化的代价，它只是最后 1.3 个 bit 的代价」「困惑度是这道门的下限，不是上限……指令遵循、工具调用、长程连贯性会先退化，而这些没有被测」「通过 1.25 只代表『这个分配是合理的』，绝不代表『模型没被损伤』」。

十一、总评

把四层验证的结果摆在一起：

所以我的判断是：这是真货，而且真得有点过分。 一个没有学术背景的独立开发者，五年、一张卡，把 27B 模型压进 12 GB，并把手写 kernel 从落后 95% 追到落后 14%——全程不碰 CUTLASS，只靠 CUDA C++、裸 mma.sync PTX 和 ncu。

但它是一个craft project（手艺项目），作者从头到尾都在强调这一点，甚至把「不参加吞吐排行榜」写进了产品定义。它真正的交付物不是 tok/s，而是一条可以从 main() 一路读到 PTX 指令的路径，以及沿途那些被诚实记录下来的错误——包括一个存活了 32 天的越界读、一个通过了所有测试的 bug、和一次「教科书算法只买来 1.2 倍」的证伪。

最后一句：在这个 AI 可以一天生成十个「高性能推理引擎 README」的年代，检验真伪的标准已经变了。不是看文档漂不漂亮，而是看它敢不敢写下自己慢多少、错在哪、以及哪些东西自己没能验证。
Mila 在这三件事上，做得比绝大多数明星项目都好。

附：本次研究的可核验清单

研究日期 2026-09-15 ｜ 方法：仓库源码与规格文档直读 + 官方博客交叉印证 + 外部工具链发布页核验 + 双路子代理并行取证。
本文所有性能数字均为作者自报，未经第三方复现。凡文档内部数字不自洽处（如 2.82 / 2.90 / 15.1 GiB），已在正文中标注，未做调和。