$ ./bin/k3 ~/k3model --trunk ~/k3trunk --preset laptop \
           --tok ~/k3model --prompt "The capital of France is" --gen 8 --incremental

--- generated text ---
 Paris.",
+            "The Eiffel
----------------------
8 tokens in 261.5 s, 32.69 s/token average
PEAK RSS for the whole run: 8.24 GB

速度很慢，但回答正确，在 8.24 GB 内存中，从 1.56 TB 的检查点运行。这是一个基础模型，所以 " Paris." 之后的内容是续写而非回复；没有 chat 模板。给它更多内存，答案不会改变，只有时间会变：

$ ./bin/k3 ~/k3model --trunk ~/k3trunk --preset server \
           --tok ~/k3model --prompt "def fibonacci(n):" --gen 28 --incremental

--- generated text ---
    if n <= 1:
        return n
    else:
        return fibonacci(n-1) + fibonacci
----------------------
28 tokens in 299.3 s, 10.69 s/token average
PEAK RSS for the whole run: 127.92 GB

本文档中的每一个数字都来自 docs/data/ 中的测量输出。

顶部是小的常驻工作集，底部是 NVMe 上的模型本身，它们之间有若干带标签的管道

密集主干（trunk）按你选择的深度保留在内存中，其余部分则流式读取；1.45 TB 的路由专家（routed experts）从不常驻，直接从其打包的 4-bit 形式进行乘法运算。结果是同一个模型可以在 8 GB 和 224 GB 内存中运行，并且在任何预算下都产生字节完全相同的输出。

关于字节存放位置的四个决策，使其从集群级部署变为笔记本级部署，而底部的答案与顶部的答案相同：

从服务器集群到普通笔记本的四个步骤，两端输出相同

第二部分 将从头开始构建两个图中的每个框，一次一个组件。

---

目录

第一部分：入门

• 要求
• 快速开始：约一分钟内完成克隆、构建和验证，无需模型
• 完整设置：生成文本的完整路径
• 用法
  • 简介
  • 提示词选项
  • 内存选项
  • 生成选项
  • 诊断选项
  • 退出码
  • 环境变量
  • 示例
• 选择预设
• 阅读运行报告
• 常见问题

第二部分：工作原理

• 问题：一个放不下的模型
• 四项缩减
• 机器及其假设
• 代码库
• 五个不变量
• 1. 从头部读取 1.56 TB 的检查点
• 2. 拒绝猜测的配置读取器
• 3. 分词器，逐字节
• 4. 缩减一：专家已经以半字节形式提供
• 5. 具有浮点契约的内核
• 6. 缩减二：KDA，内存永不增长的注意力
• 7. 缩减三：MLA，一个潜在向量代替九十六个头
• 8. 注意力残差：回望的层
• 9. 从 896 个专家中挑选 16 个
• 10. 打包主干：93 层，每层一次读取
• 11. 缩减四：流式传输主干将下限变为旋钮
• 12. 专家 LRU 缓存
• 13. 缓存应该多大？询问跟踪记录

第三部分：验证

• 门阶梯
• 先从小型预言机开始
• 在完整检查点上证明
• 第一批 token
• 持续生成：文本进，文本出

第四部分：测量

• 内存阶梯：8 GB 到 224 GB
• 未参与的缓存
• 分配胜过容量
• 测量测量本身
• 存储是全部关键
• 为什么主干不进行量化

第五部分：参考

• 范围
• 结清账目
• 文档
• 开发
• 许可证

---

第一部分：入门

要求

门槛是存储：检查点为 1.56 TB。 其他一切都是普通的。

操作系统	Linux, x86-64	使用 O_DIRECT、posix_memalign、getrusage
CPU	AVX2 + FMA	不需要 AVX-512。make portable 目标针对通用 AVX2
内存	8 GB 及以上	所有预设均可工作；内存越多越快，但结果不变
存储	约 1.7 TB 可用空间	1.56 TB 检查点 + 109 GB 打包主干，最好放在快速本地磁盘上
工具链	GCC ≥ 9 或 Clang ≥ 10	GNU make 或 CMake
Python	3.9+	用于下载、打包和分析工具；make test 不需要

分词器和配置读取器是可移植的 C99，可在任何地方构建。没有检查点，你仍然可以完成快速开始中的所有内容。

快速开始

克隆、构建并运行整个测试套件。无需检查点、无需网络、无需 Python。整个过程大约需要一分钟。

git clone https://github.com/FareedKhan-dev/kimi-k3-in-c.git
cd kimi-k3-in-c

make -j            # 几秒钟。七个 C 文件、一个编译器和 OpenMP
make test          # 不到一分钟

它要么以如下结果结束，要么失败：

GATE 1  teacher forcing : 32/32 positions match tf_pred
        generated span  : 20/20  <- must be exact
GATE 2  greedy decode   : 20/20 generated tokens match full_ids
GATE 3  incremental    : 20/20 generated tokens match full_ids  <- KV cache + carried KDA state

VERDICT: ENGINE MATCHES THE REFERENCE EXACTLY

ALL WEIGHTLESS TESTS PASSED

这就是整个引擎：每个内核、流式缓存、safetensors 读取器、配置读取器、分词器，以及一个端到端的预言机，它在一个 13 层模型上运行，该模型使用与发布版本相同的张量图构建，并与仓库中提交的 fixtures 中的 PyTorch 参考进行核对。

一个已发布的测量结果也可以当场重放，使用完整 93 层运行期间记录的跟踪数据（这需要 Python 3.9+ 和 numpy）：

python3 tools/sim_cache.py tests/fixtures/expert_trace.bin

100,096 个专家请求，重新打印 expert-cache-capacity.txt 中的容量表。

完整设置

从空目录到生成文本的六个步骤。只有第 4 步很慢。

./scripts/k3-doctor.sh 可以随时运行。它检查工具链、根据预设调整你的内存大小、测量你的存储，并打印下一步要运行的确切命令。

步骤 0. 克隆

git clone https://github.com/FareedKhan-dev/kimi-k3-in-c.git
cd kimi-k3-in-c

约 45 MB，大部分是图表和测试 fixtures。

步骤 1. 检查机器

./scripts/k3-doctor.sh

大约需要一分钟，因为它以引擎读取磁盘的方式测量你的磁盘。如果机器根本无法运行模型，它会以非零状态退出。

步骤 2. 构建

make -j

几秒钟。唯一的依赖是 C99 编译器、libm 和 OpenMP。CMake 也可以：

cmake -B build && cmake --build build -j && ctest --test-dir build

步骤 3. 在下载任何内容之前进行验证

make test

在投入 1.56 TB 下载之前，这值得一做：它证明了引擎在具有相同张量图的模型上与其参考实现匹配，并且只需要仓库本身。

步骤 4. 获取检查点

1.56 TB，所以需要数小时而非数分钟。 从 huggingface.co/settings/tokens 获取 token：

export HF_TOKEN=hf_your_token_here          # 从环境变量读取，永不回显
./scripts/download-model.sh ~/k3model       # 可恢复，重新运行以继续

脚本最后会验证分片数量、精确字节总数，然后对照已发布的数据逐一验证每个分片的大小：

verifying…
  shards : 96 (expect 96)
  bytes  : 1560936091448 (expect 1560936091448)
  shards : all 96 match their published sizes individually
  RESULT : byte-exact match

部分下载不会大声失败；它会产生错误的 token。将这里的 FAIL 视为停止信号。逐分片检查也将“重新下载 1.56 TB”变成“重新下载这个 17 GB 的文件”，并且它能捕获总数无法捕获的一种情况：两个分片以相反方向错误了相同的量。

步骤 5. 打包主干

./scripts/pack-trunk.sh ~/k3model ~/k3trunk

大约四分钟，只需一次。它将 93 个密集层重写为一个 109 GB 的文件，其中层 L 位于已知偏移量处，可以通过一次调用读取。这就是将内存需求变成旋钮的关键。 将输出放在你最快的磁盘上。

步骤 6. 运行

./bin/k3 ~/k3model --trunk ~/k3trunk --preset workstation \
         --tok ~/k3model --prompt "The capital of France is" --gen 8 --incremental

分词器随检查点一起提供，这就是为什么 --tok 指向模型目录。

所有内容的最终位置

kimi-k3-in-c/    ~45 MB   源码、文档、图片和 bin/k3
~/k3model/      1.56 TB   96 个分片 · config.json · tiktoken.model · tokenizer_config.json
~/k3trunk/       109 GB   trunk.bin · trunk.json，放在你最快的磁盘上

任何一次运行的第一个 token 都会从磁盘加载所有固定层，在 server 预设下约为 108 GB，因此所需时间远长于稳定速率。该成本每次运行支付一次，而不是每个 token 支付一次。

用法

简介

k3 <model_dir> [prompt] [memory] [generation] [diagnostics]

<model_dir> 是存放 .safetensors 分片的目录。任何运行都需要它，但 --help、--version 和 --list-presets 不需要它：

./bin/k3 --help
./bin/k3 --version
./bin/k3 --list-presets

提示词选项

以下选项必须且只能提供一个。不传或传多个都是用法错误（退出码 2）。

标志	参数
--prompt	TEXT	对 TEXT 进行分词并运行。需要 --tok。
--prompt-file	PATH	对文件的字节进行分词。需要 --tok。 对于任何非 ASCII 内容，这是首选：shell 会重新编码 argv，而文件是逐字读取的
--ids	1,2,3	直接提供 token id。完全不加载分词器，因此可以在没有分词器文件的机器上工作。这是测试使用的可复现通道

# 文本输入
./bin/k3 ~/k3model --tok ~/k3model --prompt "The capital of France is" ...

# 文本输入，来自文件。用于中文、日文、韩文、表情符号、重音符号
printf 'La capitale de la France est' > /tmp/p.txt
./bin/k3 ~/k3model --tok ~/k3model --prompt-file /tmp/p.txt ...

# id 输入，id 输出，无需分词器
./bin/k3 ~/k3model --ids 1008,10484,318,15383,387 ...

内存选项

标志	参数	默认值
--preset	NAME	无	laptop · desktop · workstation · server · max。设置以下两个预算
--trunk	DIR	关闭	来自步骤 5 的打包主干目录。这是启用流式传输的关键。 没有它，主干将完全常驻加载，约 113.5 GB
--trunk-gb	X	16	固定层加流式环形缓冲区的预算
--cache-gb	X	64	路由专家 LRU 缓存的预算

--preset 和两个 -gb 标志设置相同的两个数字，因此预设只是简写。如果混合使用，顺序很重要：后面的标志优先，所以 --preset server --cache-gb 40 会给你 server 的主干预算和 40 GB 的缓存。

没有 --trunk 的 --preset 没有实际作用。 每个预设都假设主干是流式传输的。省略 --trunk，引擎会加载全部 113.5 GB 常驻内存，无论你要求的预算是多少。

生成选项

标志	参数	默认值
--gen	N	8	要生成的 token 数。上限 4096；提示词最多 32768 个 token
--incremental	无	关闭	在 token 之间携带 KV 缓存和循环状态，而不是重新运行整个前缀
--tok	DIR	无	存放 tiktoken.model 和 tokenizer_config.json 的目录

任何长度的生成都必须传递 --incremental。 没有它，每一步都会重新运行整个前缀，这是 O(T²)；有了它，第 0 步支付提示词的成本，之后的每一步成本固定。两条路径都经过验证，必须产生相同的 token，因此这是一个纯粹的速度选择。

诊断选项

标志	参数
--config	PATH	模型配置；默认为 <model_dir>/config.json
--layers	N	仅绑定前 N 层，用于部分分片集
--out	FILE	JSON 结果（默认 k3_run.json）
--dump-logits	PATH	第一步的 float32 logits，用于逐元素比较
--dump-cache-trace	DIR	写入 expert_hist.json 和 expert_trace.bin，tools/sim_cache.py 可重放

退出码

脚本可以依赖这些。

0	成功
1	张量绑定失败，或前向传播失败
2	用法错误，或无法自信读取的配置；引擎拒绝猜测
4	运行完成，但至少一个路由专家加载失败，因此发出的 id 不可靠。与 1 不同，因为进程其余部分成功，并且这是捕获静默数值损坏的代码

环境变量

变量	使用者
HF_TOKEN	download-model.sh	HuggingFace token，从环境变量读取，永不回显
OMP_NUM_THREADS	引擎	线程数，默认为所有核心
K3_TOK_FILES	分词器工具和 CI	存放 tiktoken.model 的目录，当它不在默认位置时
K3_MODEL_DIR	tools/budget.py	检查点目录，当未作为参数给出时

示例

# 最小可能的运行，8 GB 下限。
./bin/k3 ~/k3model --trunk ~/k3trunk --preset laptop \
         --tok ~/k3model --prompt "Hello! My name is" --gen 16 --incremental

# 每 GB 最快。固定 93 个主干层中的 90 个。
./bin/k3 ~/k3model --trunk ~/k3trunk --preset server \
         --tok ~/k3model --prompt "def fibonacci(n):" --gen 28 --incremental

# 手动调整拆分而不是预设：所有内存都给主干。
./bin/k3 ~/k3model --trunk ~/k3trunk --trunk-gb 110 --cache-gb 13 \
         --tok ~/k3model --prompt-file prompt.txt --gen 32 --incremental

# 可复现：id 输入，id 输出，无分词器，JSON 结果。
./bin/k3 ~/k3model --trunk ~/k3trunk --preset desktop \
         --ids 1008,10484,318,15383,387 --gen 8 --incremental --out run.json

# 捕获缓存跟踪，然后在任何容量下离线重放。
./bin/k3 ~/k3model --trunk ~/k3trunk --preset workstation \
         --ids 1008,10484,318,15383,387 --gen 8 --incremental \
         --dump-cache-trace /tmp/trace
python3 tools/sim_cache.py /tmp/trace/expert_trace.bin

# 与 PyTorch 参考进行逐元素 logit 比较。
./bin/k3 ~/k3model --trunk ~/k3trunk --preset server \
         --ids 3,4,5,6,7 --gen 1 --dump-logits /tmp/c_logits.bin
python3 tools/cmp_logits.py /tmp/c_logits.bin ref_logits.json

# 部分分片集：仅绑定前 8 层。
./bin/k3 ~/k3model --trunk ~/k3trunk --layers 8 \
         --ids 1,2,3 --gen 1

# 在硬性内存上限下运行，这是阶梯测量的方式。
systemd-run --scope --user -q -p MemoryMax=8G -p MemorySwapMax=0 \
  ./bin/k3 ~/k3model --trunk ~/k3trunk --trunk-gb 2.5 --cache-gb 0.5 \
           --ids 1008,10484,318,15383,387 --gen 8 --incremental

选择预设

$ ./bin/k3 --list-presets
presets (trunk / expert-cache, in GB):
  laptop          3.0 / 1.0     8.2 GB peak RSS. The floor. Runs, slowly.
  desktop        16.0 / 10.0    31.9 GB peak RSS.
  workstation    60.0 / 30.0    95.5 GB peak RSS; the expert cache starts to matter here.
  server        110.0 / 13.0    ~128 GB peak RSS; 90 of 93 trunk layers pinned. Fastest.
  max           110.0 / 109.0   ~224 GB peak RSS; trunk pinned and a large expert cache.

All presets stream the trunk, so they need --trunk <packed_dir>.
Run scripts/k3-doctor.sh to see which one this machine fits.

每个预设实际消耗的内存

边界来自实测阶梯，doctor 脚本根据 MemAvailable 而不是 MemTotal 来判断：

if   [ "$AVAIL_GB" -ge 192 ]; then PRESET=server;      EXPECT="~19-21 s/token"
elif [ "$AVAIL_GB" -ge  96 ]; then PRESET=workstation; EXPECT="~24 s/token"
elif [ "$AVAIL_GB" -ge  32 ]; then PRESET=desktop;     EXPECT="~28-31 s/token"
elif [ "$AVAIL_GB" -ge  10 ]; then PRESET=laptop;      EXPECT="~32 s/token"
else PRESET=""; fi

在选择之前，有两件事值得了解：

• 在这些测量中，max 并不比 server 快。 额外的 96 GB 在噪声底线之外没有带来任何收益。
• 在给专家缓存分配内存之前，先给主干分配内存。 在固定的 128 GB 预算下，这带来了 1.69 倍的提升。分配胜过容量 有相关数据。

阅读运行报告

引擎打印内存计划，然后每个生成的 token 打印一行，最后打印摘要。以下摘自一次 workstation 运行：

cache [final step]
  requests     : 1472  hits 1472 (100.00%)  misses 0  evictions 729
                 TRUE resident hit rate 50.48%
I/O share of wall clock: 71.1%  (trunk 62.4 s + experts 34.2 s of 135.8 s)
trunk [final]
  pinned 48/93 layers, ring 1 slots
  read 368.65 GB in 62.40 s (5908 MB/s)
PEAK RSS for the whole run: 94.74 GB   <- quote this, not the plan

三个数字承载了关键信息：

• TRUE resident hit rate：从 RAM 提供的专家。原始的 hits 计数器也计算了预取器刚刚从磁盘拉取的专家，因此它在任何缓存大小下都显示 100%；常驻命中率打印在其下方。
• I/O share of wall clock：整个运行的磁盘时间占总时间的比例，在阶梯上测得在 41% 到 61% 之间。
• PEAK RSS：来自 getrusage，在运行之后。这是内存数字；预先的计划略高于此值。

常见问题

即使在小预设下，内存也接近 113 GB。 省略了 --trunk。没有打包的主干目录，整个主干都会常驻加载；每个预设都假设流式传输。

非 ASCII 提示词分词异常。 Shell 会重新编码 argv，因此引擎接收到的字节与你输入的不同。将提示词放入文件并使用 --prompt-file，它会逐字读取。

--prompt/--prompt-file 需要 --tok DIR。 分词器随检查点提供，所以添加 --tok ~/k3model。引擎会退出而不是猜测词汇表的位置。要完全跳过分词器，请使用 --ids 传递 token id。

吞吐量远低于表格。 几乎总是存储问题。python3 tools/devbw.py <file-on-that-disk> 以引擎读取磁盘的方式测量磁盘，使用大随机 O_DIRECT 读取，队列深度为 1 和 16，这是 dd 做不到的。网络卷比本地 NVMe 慢几倍；将 ~/k3trunk 保持在本地。

运行因 KV 缓存而拒绝启动。 无论预算如何，上下文每位置大约需要 2.37 MB，引擎会预先计算这一点，而不是在一小时后才发现。缩短请求，或去掉 --incremental，后者完全不携带 KV 缓存。

是否需要完整的 1.56 TB？ 对于生成，是的。对于开发，不需要：make test 完全不需要任何东西，--layers N 可以针对部分分片集运行。

macOS、Windows、WSL？ 引擎针对 Linux。分词器和配置读取器是可移植的 C99，并在 CI 中以可移植方式构建。

---

第二部分：工作原理

问题：一个放不下的模型

Kimi K3 拥有 2.78 万亿参数，发布时大小为 1.56 TB。没有消费级机器能容纳它，等待更好的硬件也无济于事，因为瓶颈不是速度，而是容量。

但它是一个专家混合模型，所以对于任何给定的 token，每层 896 个专家中只有 16 个被激活，其余的在磁盘上休眠。将始终在线的部分保留在内存中，流式传输休眠的专家，它就能在 8.24 GB 内存中、单 CPU、无 GPU 的情况下运行。

朴素的需求是每个参数数量所暗示的那个。

朴素需求：每个参数都以 bf16 常驻

所以 5.56 TB 是需要击败的数字。

一个 token 唤醒 16 个专家，880 个保持休眠

Kimi K3 有 93 层。第 0 层是普通的密集前馈层，所以其他 92 层进行路由，每层从 896 个专家中选出前 16 个。

每层只有 896 个专家中的 16 个被激活，所以模型的大部分处于休眠状态

对于任何给定的 token，大约有 1040 亿参数处于激活状态，占 2.78 万亿的 3.7%。其余 96.3% 仍然需要存在于某个可访问的地方，但不必在 RAM 中。

计算磁盘上的实际字节数，而不是猜测：

=== shard census: what the 1.56 TB actually is ===
shards            : 96
total bytes       : 1560936091448  (1.56 TB)

--- routed experts (the part that is streamed, never resident) ---
  experts total     : 82,432   (896 routed x 92 MoE layers)
  bytes per expert  : 17,547,264  exactly
                      = 33,030,144 params x 0.53125 bytes
                      = 0.5 bytes/nibble + 1/32 byte for the shared E8M0 scale
  routed expert set : 82,432 x 17,547,264 = 1.447 TB

有 82,432 个路由专家，每个恰好占用 17,547,264 字节。它们合计为 1.447 TB，占整个检查点的 93%。其他所有内容（注意力投影、路由器、归一化、嵌入）是剩余的 7%。

1.56 TB 的分布：93% 是永不加载的专家

这个普查就是一张图里的整个策略。如果这 1.447 TB 可以可访问但从不常驻，那么在编写任何内核之前，内存问题就会缩小一个数量级以上。

始终激活的集合：bf16 下 113.49 GB，其他一切均可流式传输

剩下的是 56,743,648,000 个参数，即 bfloat16 下的 113.49 GB。其中，108.81 GB 是每层的密集主干，4.70 GB 是嵌入表加上输出头。

四项缩减

• 5,560 GB：每个参数都是 bfloat16，这是我们开始的地方。
• 1,560 GB：检查点发布时的大小，因为专家已经以每个权重半字节的形式提供。
• 113.49 GB：一旦路由意味着专家从不加载，必须常驻的内容。
• 8.24 GB：一旦主干被流式传输而不是持有，实测的大小。

四项缩减，两端输出相同

端到端，这是相对于 bfloat16 模型的 675 倍缩减，相对于发布检查点的 189 倍缩减。没有近似，没有丢弃任何权重：阶梯底部的输出与顶部的输出逐字节相同。本文档顶部的图表就是按比例绘制的该账目。

机器及其假设

这里的每个测量都来自一台工作站：双路 AMD EPYC 7763，124 核，无 SMT，228 GB 内存，3.2 TB NVMe。它还有四块 NVIDIA L40 GPU，在整个过程中完全闲置，因为这个引擎没有 GPU 路径。

--- ISA (note: AVX2 present, AVX-512 ABSENT) ---
avx avx2 fma sse4_2

--- memory ---
Mem:           228Gi       5.1Gi       207Gi       3.1Mi        18Gi       223Gi
MemTotal:       239308464 kB
MemAvailable:   233961008 kB
Hugepagesize:       2048 kB

没有 AVX-512。 引擎只需要 AVX2 和 FMA，这是过去十年任何台式机 CPU 都有的指令集。

存储数字比 CPU 数字更重要，而且有一个与预期相反。

--- storage bandwidth, measured ---
O_DIRECT cold : 3.2 GB/s     (dd bs=4M iflag=direct after drop_caches)
buffered warm : 2.3 GB/s
engine, trunk : 5373-6064 MB/s sustained during runs
NOTE O_DIRECT is FASTER than buffered here. That is the opposite of the usual
expectation, and it is why the engine opens the trunk O_DIRECT.

使用 O_DIRECT 读取，完全绕过页缓存，在这里比通过它读取更快。这一项测量决定了整个 I/O 设计。

一个二进制，四种机器，一个相同的答案

一项卫生措施，因为负载高的机器很容易测量不准：

--- measurement hygiene ---
unattended-upgrades: STOPPED and DISABLED before measurement (was using ~63% of a
  core during the smoke run).
apt-daily.timer and apt-daily-upgrade.timer: DISABLED

后台的软件包更新程序占用大部分核心，对计时的影响比大多数优化都大，所以在测量任何东西之前它必须关闭。

引擎实际需要多少内存？手动乘以配置值会以一种有启发性的方式给出错误答案，这就是 tools/budget.py 存在的原因：

# Streamable only if ROUTED. The 2 SHARED experts sit in the same namespace and
# are NOT streamable, which is where hand arithmetic goes wrong.
def classify(name: str) -> str:
    if ".block_sparse_moe.experts." in name:
        return "routed_expert"          # streamable: only 16 of 896 per token
    if ".block_sparse_moe.shared_expert" in name:
        return "shared_expert"          # RESIDENT: runs on every token
    if ".self_attn." in name:
        return "attention"              # resident
    if "embed_tokens" in name or "lm_head" in name:
        return "embedding"              # resident
    return "other"                      # norms, router gates, biases: resident

两个共享专家在每个 token 上运行，所以它们属于常驻集，即使它们的张量名与路由专家相邻。弄错这一点会使下限看起来比实际更小，这是最糟糕的错误方向。

---

代码库

六个 C 文件编译成一个二进制。没有 BLAS，没有 PyTorch，没有 ONNX runtime，没有 GPU 库。唯一的依赖是 libm 和 OpenMP。

include/k3/
  k3.h              # the public header: config, weights, every kernel prototype
  k3_cfg.h          # config reader, header-only, refuses to substitute defaults
src/
  core/k3_ops.c     # every numeric kernel: RMSNorm, KDA, MLA, MoE, MXFP4 matmul
  io/k3_st.c        # safetensors reader, hand-written JSON scan, O_DIRECT reads
  io/k3_load.c      # locating one expert's bytes inside a shard
  io/k3_trunk.c     # streaming the dense trunk, pinned prefix plus a ring slot
  cache/k3_cache.c  # the routed-expert LRU cache and its batch prefetch
  model/k3_bind.c   # binding checkpoint tensor names to kernel arguments
  tokenizer/k3_tok.h# byte-level BPE loaded from tiktoken.model
  cli/k3_run.c      # the k3 binary: memory plan, decode loop, reporting
tools/              # python: pack the trunk, replay the cache, verify against torch
benchmarks/         # the cgroup memory ladder and the split sweep
tests/              # fixtures, the tiny oracle, the 93-layer conformance run

CFLAGS = -O3 -std=gnu99 -Wall -Wextra -Wpointer-arith -Wshadow -Wvla \
         -march=native -fopenmp -ffp-contract=off
LDFLAGS = -lm -fopenmp

看起来不寻常的标志是 -ffp-contract=off。默认情况下，编译器可能会将乘法和加法融合成单个 FMA，这会改变舍入。这通常是好事。但在这里是个问题，因为标量路径、OpenMP 路径和 AVX2 路径必须产生位相同的结果，这样性能变化永远不会悄悄变成精度变化。

一个 C 文件加上小头文件变成一个微小的静态二进制

1. build, warnings are failures
  -> clean build, no diagnostics
  test_ops          97784 bytes
  k3_model          89392 bytes
  k3_run           179736 bytes

整个推理引擎是 179,736 字节，一个 176 KB 的二进制文件，其工作是运行一个 1.56 TB 的模型。

一个运行 1.56 TB 模型的 176 KB 二进制文件

一项检查跨越了机器。分词器在 Windows 和 Linux 下对同一输入文件运行：

=== cross-platform tokenizer determinism ===
  Linux   gcc 13.3.0  x86_64
  Windows gcc 16.1.0  x86_64
  input   src/k3.h (24,499 bytes) -> 6,862 ids
  result  IDENTICAL id streams

  (a naive md5 of stdout DIFFERS by one byte: Windows text-mode stdout writes the
   trailing newline as CRLF. That is the pipe, not the tokenizer.)

两个操作系统上的两个编译器从相同的 24,499 字节产生相同的 6,862 个 token id。md5 总和恰好相差一个字节，原因是 shell 添加的行尾，而不是分词器做的任何事情。

五个不变量

公共头文件以五个必须成立的不变量开头。每一个都是一个看似合理的实现会产生一个能运行、能输出流畅文本但却是错误的模型的地方，没有崩溃，没有 NaN 来警告你。

1. A_log 按头索引，而不是按通道索引。 检查点提供 head_dim 个浮点数，但只有前 num_heads 个是有意义的；其余是填充。
2. UT 变换的逆是 (I + Akk)^-1。 符号不是约定。
3. Aqk 保留其对角线，而 Akk 不保留。
4. MLA 使用 NoPE，但 64 个 rope 维度仍然存在并且仍然被缓存。 只有旋转被省略；删除这些槽会改变头宽度。
5. MoE 路由偏置仅引导选择。 组合权重来自无偏 sigmoid 分数。

下面每个组件出现时都会勾选对应项。

1. 从头部读取 1.56 TB 的检查点

检查点是 96 个 safetensors 文件。该格式故意设计得很简单，这使得将 1.56 TB 视为索引而不是数据成为可能。

safetensors：一个长度，一个头部，然后是已知偏移量处的原始字节

每个文件以 8 字节小端长度开头，然后是那么多字节的 JSON 描述每个张量，然后是背靠背的原始张量字节。没有压缩，没有交错。

索引分片，按需读取确切字节，然后丢弃页面

不使用 JSON 库。头部可能有几十 MB，每个张量只需要四个字段，所以读取器直接扫描它。

/* Walk the header once, copy nothing we do not need. `p` sits just past the
 * opening quote of the tensor name. */
static const char *st_scan_entry(const char *p, const char *end, K3Tensor *t)
{
    const char *q = memchr(p, '"', (size_t)(end - p));
    if (!q || (size_t)(q - p) >= sizeof t->name) return NULL;
    memcpy(t->name, p, (size_t)(q - p));
    t->name[q - p] = '\0';

    const char *d = st_find_key(q, end, "dtype");
    if (!d) return NULL;
    t->dtype = st_dtype_code(d);

    const char *s = st_find_key(q, end, "shape");
    if (!s) return NULL;
    t->rank = 0;
    t->nelem = 1;
    for (const char *c = s; c < end && *c != ']'; c++) {
        if (*c >= '0' && *c <= '9') {
            long v = strtol(c, (char **)&c, 10);
            if (t->rank >= K3_ST_MAXRANK) return NULL;
            t->shape[t->rank++] = v;
            t->nelem *= (size_t)v;
        }
    }

    /* offsets are RELATIVE to the start of the data section */
    const char *o = st_find_key(q, end, "data_offsets");
    if (!o) return NULL;
    t->off  = (size_t)strtoull(o, (char **)&o, 10);
    while (o < end && (*o < '0' || *o > '9')) o++;
    t->nbytes = (size_t)strtoull(o, (char **)&o, 10) - t->off;
    return o;
}

每个张量都进入一个以名称哈希为键的哈希表。哈希的选择不是任意的。

/* Names are long and share deep prefixes
 * ("language_model.model.layers.N.block_sparse_moe.experts.M...."), so the hash must
 * mix every byte; a prefix-only or length-only hash would pile every expert of a
 * layer into one bucket. */
static uint64_t fnv1a(const char *s)
{
    uint64_t h = 1469598103934665603ull;
    while (*s) { h ^= (unsigned char)*s++; h *= 1099511628211ull; }
    return h;
}

五十万个张量名都以相同的四十个字符开头，这对哈希函数来说是一个真正敌对的输入。FNV-1a 在每个字节上混合，所以名称末尾的专家索引仍然会影响结果。

之后读取张量有一个小问题：O_DIRECT 要求偏移量和长度是块大小的倍数，而张量在前一个张量结束的地方开始。

int64_t k3_st_read_aligned(const K3St *s, int shard, int64_t off, int64_t nbytes,
                           void *buf, int64_t bufcap, int64_t *payload_off)
{
    /* widen outward to the enclosing aligned window */
    const int64_t lo  = off & ~(int64_t)(K3_ST_ALIGN - 1);
    const int64_t hi  = (off + nbytes + K3_ST_ALIGN - 1) & ~(int64_t)(K3_ST_ALIGN - 1);
    const int64_t len = hi - lo;
    const int64_t pad = off - lo;
    if (len > bufcap) return 0;
    if (payload_off) *payload_off = pad;

    int64_t got = 0;
    while (got < len) {
        ssize_t r = pread(dfd, (char *)buf + got, (size_t)(len - got), (off_t)(lo + got));
        if (r <= 0) break;      /* the last window may run past EOF */
        got += r;
    }
    return got >= pad + nbytes ? nbytes : (got > pad ? got - pad : 0);
}

注意短读时是 break 而不是失败：分片的最后一个对齐窗口会延伸到文件末尾之外，这是预期的，所以返回值检查的是有效载荷是否被覆盖，而不是整个窗口是否被覆盖。

indexed 497220 tensors from 96 shards in 0.27 s

大约四分之一秒内索引了五十万个张量。 这是让之后一切成为可能的关键：引擎从不读取它不需要的分片，所以磁盘上的 1.56 TB 是一个目录，而不是一个工作集。

一个与自己一致的解析器证明不了什么，所以索引被转储并在 Python 中独立重新解析，比较 dtype、形状、偏移量和扩展后的浮点位模式。

# Bit patterns, not tolerances: widening bf16 to f32 is lossless.
c_bits = np.asarray(c_values[name], dtype=np.float32).view(np.uint32)
p_bits = ref.astype(np.float32).view(np.uint32)
if not np.array_equal(c_bits, p_bits):
    bad = int(np.count_nonzero(c_bits != p_bits))
    fail(f"{name}: {bad} of {c_bits.size} float32 bit patterns differ")

=== shard verification ===
shards: 96
bytes:  1560936091448
expected: 1560936091448
RESULT: EXACT MATCH

96 个分片，1,560,936,091,448 字节，逐个文件验证

2. 拒绝猜测的配置读取器

模型的维度来自检查点自己的 config.json，这是不变量四可能悄悄咬人的第一个地方。

基于一的层索引，92 和 93 都是 MLA，这是设计使然

Kimi K3 交替使用两种注意力机制。大多数层使用一种，每四层使用另一种，最后两层都是第二种，所以最后一层总是做全局注意力。配置明确列出了这些层，并且列表是基于一的。

--- every value below is READ from the checkpoint, not assumed ---
config: config.json (nested shape) | hidden=7168 layers=93 vocab=163840
        | 24 MLA + 69 KDA | experts 896 top16 shared2 | latent=3584

--- KDA/MLA layer map (ONE-based, from full_attn_layers) ---
full_attn_layers (24, all MLA): 4,8,12,16,20,24,28,32,36,40,44,48,52,56,60,64,68,72,76,80,84,88,92,93
  note 92 AND 93 are both MLA - the report (2.1) places an extra Gated MLA layer
  at the end of the backbone so the final layer always does global attention.
kda_layers (69): every other layer.

这些数字中的每一个都是从文件中读取的；没有一个是编译进去的。它们落在一个结构体中，这就是一屏上的整个模型：

typedef struct {
    int hidden;            /* 7168  */
    int n_layers;          /* 93    */
    int vocab;             /* 163840 */
    float rms_eps;         /* 1e-5  */

    /* Kimi Delta Attention. 69 of the 93 layers. */
    int kda_heads;         /* 96    */
    int kda_head_dim;      /* 128, and d_k == d_v */
    int conv_k;            /* 4, depthwise, causal, SiLU fused */
    float gate_lb;         /* -5.0, the decay lower bound */

    /* Gated MLA. 24 of the 93 layers. */
    int n_heads;           /* 96    */
    int q_lora;            /* 1536  */
    int kv_lora;           /* 512   */
    int qk_nope;           /* 128   */
    int qk_rope;           /* 64, PRESENT BUT NEVER ROTATED */
    int v_head;            /* 128   */
    int mla_out_gate;      /* 1     */

    /* Stable LatentMoE. 92 of the 93 layers. */
    int n_experts;         /* 896   */
    int topk;              /* 16    */
    int n_shared;          /* 2, full width, added UNWEIGHTED */
    int latent;            /* 3584, the routed-expert width */
    int moe_inter;         /* 3072  */
    float routed_scale;    /* 1.0   */
    int moe_renorm;        /* 1     */
    int latent_norm;       /* 1, RMSNorm on the AGGREGATE, not per expert */

    /* the single dense layer, layer 0 */
    int first_dense;       /* 1     */
    int dense_inter;       /* 33792 */

    int attn_res_block;    /* 12. Boundaries fire when layer_idx % this == 0. */
    float situ_b1;         /* 4.0   */
    float situ_b2;         /* 25.0  */

    int  n_full_attn;      /* 24 */
    int *full_attn;        /* ONE-BASED layer indices */
} K3Cfg;

这个结构体是检查点和每个内核之间的契约。如果它是对的，模型就是 Kimi K3。如果任何字段错误，模型就是另一个仍然会说英语的东西。

拒绝而不是猜测，因为猜测的字段会给你一个不同的模型

考虑一个宽容的读取器会做什么。发布的配置将其字段嵌套得比 fixture 深一层，所以一个只知道扁平形状的读取器找不到它认识的东西。如果它然后填充默认值，会发生两件事：SiTU beta 得到 4.0 和 25.0，这是正确的值，所以看起来没有错，而 full_attn_layers 返回空，所以所有 93 层都作为 KDA 运行，24 个全局注意力层消失了。模型加载、流式传输、解码，并从一个不是 Kimi K3 的架构产生语法正确的英语。

/* An absent field is an ERROR, never a default. Missing names are accumulated so
 * the message lists all of them at once. */
static int cfg_req_int(jval root, const char *key, int *out,
                       const char **missing, int *nmissing)
{
    jval v = json_get(root, key);
    if (v.type != JSON_NUM) {                 /* absent OR the wrong type */
        if (*nmissing < K3_CFG_MAXMISS) missing[(*nmissing)++] = key;
        return 0;
    }
    *out = (int)v.num;
    return 1;
}

  [no_layermap]
    k3_cfg: no_layermap.json is missing 1 required field(s):
        full_attn_layers
      refusing to substitute defaults: a config this reader cannot
      fully understand would silently produce a DIFFERENT model.
      ok    correctly rejected no_layermap.json

  [bad_layer_index]
    k3_cfg: bad_layer_index.json full_attn_layers[2] = 999 is outside 1..93
        (the list is ONE-based)
      ok    correctly rejected bad_layer_index.json

配置读取器大约有一百五十行，是项目中最无聊的代码，它是能给你一个不同模型而不告诉你的两个地方之一。

3. 分词器，逐字节

另一个是分词器。Kimi K3 使用字节级 BPE，有 163,584 个 rank 加上 256 个特殊 token，以 tiktoken.model 文件形式提供。

每个案例都通过文件，而不是 argv

加载器将该文件直接读入 vendored BPE 结构。它基于三个假设，每个假设都会产生一个在 ASCII 上完美工作而在其他一切上分叉的分词器：

• 合并键是字节，而不是码点。 碰巧解码为有效 UTF-8 的键仍然必须被视为其原始字节。
• Rank 来自文件。 它们不是在加载时从频率推导出来的。
• 添加的 token 块附加在 rank 之后，所以添加的 token 的 id 是 163,584 加上其索引，而不是它在合并表中的位置。

测试将 C 分词器与 Python tiktoken 库逐案例比较，通过文件而不是命令行参数：

oracle   : tiktoken 0.13.0
method   : token-for-token comparison; every case passed through a FILE, never argv
           (argv is re-encoded to the active code page on Windows and would compare
            different bytes on every non-ASCII case)

  PASS  han only                 2 ids
  PASS  japanese                 6 ids
  PASS  korean                   5 ids
  PASS  cyrillic                 4 ids
  PASS  arabic                   7 ids
  PASS  emoji zwj                5 ids
  PASS  code python             11 ids
  PASS  json                    19 ids

tokenizer parity: 45/45 cases match

然后整个文件被推过并解码回来：

roundtrip: 48353 bytes -> 14797 ids -> 48353 bytes : PASS   <- k3_ops.c
roundtrip: 24499 bytes -> 6862 ids -> 24499 bytes : PASS   <- k3.h
roundtrip: 201775 bytes -> 52671 ids -> 201775 bytes : PASS   <- REPORT.md
roundtrip: 53444 bytes -> 12145 ids -> 53444 bytes : PASS   <- modeling_kimi_k3.py

四个文件进，字节相同的文件出

两百 KB 的 markdown 变成 52,671 个 token id，并作为完全相同的两百 KB 返回。后面关于相同输出的每一个声明都依赖于分词器的确定性。

/* Greedily merge the lowest-rank adjacent pair. Everything here is BYTES. */
static int tok_encode_piece(const Tok *t, const unsigned char *p, int n, int *out)
{
    int parts[K3_TOK_MAXPIECE + 1], np = n + 1;
    for (int i = 0; i <= n; i++) parts[i] = i;          /* byte boundaries */

    for (;;) {
        int best = -1, bestrank = INT_MAX;
        for (int i = 0; i + 2 < np; i++) {
            const int r = tok_rank(t, p + parts[i], parts[i + 2] - parts[i]);
            if (r >= 0 && r < bestrank) { bestrank = r; best = i; }
        }
        if (best < 0) break;                            /* no mergeable pair left */
        memmove(&parts[best + 1], &parts[best + 2],
                (size_t)(np - best - 2) * sizeof(int));
        np--;
    }

    for (int i = 0; i + 1 < np; i++)
        out[i] = tok_rank(t, p + parts[i], parts[i + 1] - parts[i]);
    return np - 1;
}

循环维护一个切片边界列表，并反复合并 rank 最低的相邻对，这使结果独立于任何平局打破顺序。

4. 缩减一：专家已经以半字节形式提供

四项缩减中的第一项，也是最大的一项。

路由专家不是以 bfloat16 提供的。它们以 MXFP4 提供，这是一种微缩放 4 位浮点格式。每个权重是一个 4 位半字节，索引一个 16 项表，每 32 个连续权重共享一个 8 位指数。

MXFP4：一个 4 位半字节，每 32 个权重共享一个 8 位指数

一个字节携带两个权重，低半字节是偶数

每个权重半字节加上共享缩放，正好给出一个专家

半字节加上三十二分之一字节是每个权重 0.53125 字节，一个专家有 33,030,144 个参数，所以一个专家正好是 17,547,264 字节。这与普查结果逐字节匹配。

该格式是针对发布的检查点而不是文档进行验证的：

{
  "note": "Kimi K3 MXFP4 bytes from the released checkpoint. w = E2M1[nibble] * 2^(scale - 127), one scale per 32 elements.",
  "source": "language_model.model.layers.1.block_sparse_moe.experts.0.w1",
  "rows": 64, "packed_cols": 1792, "scale_cols": 112,
  "logical_width": 3584, "group_size": 32,
  "e2m1_lut": [0.0, 0.5, 1.0, 1.5, 2.0, 3.0, 4.0, 6.0,
              -0.0, -0.5, -1.0, -1.5, -2.0, -3.0, -4.0, -6.0]
}

解码的两半都是查找表，构建它们是该格式唯一需要的设置。

/* E2M1: sign, two exponent bits, one mantissa bit. Sixteen values in total. */
static const float K3_E2M1[16] = {
    0.0f,  0.5f,  1.0f,  1.5f,  2.0f,  3.0f,  4.0f,  6.0f,
   -0.0f, -0.5f, -1.0f, -1.5f, -2.0f, -3.0f, -4.0f, -6.0f
};

/* Byte -> its two weights, so the inner loop does one lookup, not two shifts. */
static void k3_pair_init(void)
{
    for (int b = 0; b < 256; b++) {
        K3_E2M1_PAIR[b][0] = K3_E2M1[b & 0x0F];   /* low nibble  = EVEN element */
        K3_E2M1_PAIR[b][1] = K3_E2M1[b >> 4];     /* high nibble = ODD  element */
    }
}

/* Scale byte -> power of two. 255 is NaN by spec, mapped to 0 to contain damage. */
static void k3_e8m0_init(void)
{
    for (int b = 0; b < 256; b++)
        K3_E8M0[b] = (b == 255) ? 0.0f : ldexpf(1.0f, b - 127);
}

现在是半字节顺序，这是一个没有统计量可以检查的约定。

低半字节是偶数元素，反转它是静默错误的

"expected_swapped_nibbles": {
  "note": "what you get if the low nibble is treated as the ODD element.
           Statistics are identical; positions are wrong."
}

交换版本的每个均值、每个标准差、每个直方图都与正确版本相同，因为它是相同的数字多重集。只有位置不同。检查分布的验证会通过一个每对相邻权重都转置的矩阵。

使用这些权重的明显方法是将它们解码为浮点数，然后进行正常的矩阵乘法。计算一下成本：

反量化会花费什么，这就是为什么我们从半字节直接相乘

一个 17.55 MB 的专家扩展到 float32 后变成 132 MB。每个 token 在 92 层中接触 16 个专家（1,472 个专家），所以解码它们全部意味着在任何一个乘加之前，要写出每 token 194 GB 的纯格式转换。

每个权重半字节节省 4 TB，永不反量化每 token 节省 194 GB

所以没有任何东西被反量化。矩阵乘法直接读取打包的半字节。

/* y[rows] = x[in] . W[rows][in], W stored as MXFP4. Nothing is dequantized. */
void k3_matmul_mxfp4(float *y, const float *x, const unsigned char *packed,
                     const unsigned char *scales, int in, int rows, int group)
{
    const int ngroup = (in + group - 1) / group;
    const int rowbytes = (in + 1) / 2;              /* two nibbles per byte */

#pragma omp parallel for schedule(static)
    for (int o = 0; o < rows; o++) {
        const unsigned char *pb = packed + (size_t)o * rowbytes;
        const unsigned char *sb = scales + (size_t)o * ngroup;
        double acc = 0.0;                            /* double, always */

        for (int g = 0; g < ngroup; g++) {
            const float s = K3_E8M0[sb[g]];
            if (s == 0.0f) continue;                 /* a NaN scale zeroes the group */

            const int i0 = g * group;
            const int n = (i0 + group <= in) ? group : (in - i0);
            float wf[64];

            /* low nibble is the EVEN weight, high nibble the ODD one */
            const int half = n / 2;
            for (int k = 0; k < half; k++) {
                const unsigned char b = pb[(i0 / 2) + k];
                wf[2 * k]     = K3_E2M1_PAIR[b][0];  /* low  nibble -> even index */
                wf[2 * k + 1] = K3_E2M1_PAIR[b][1];  /* high nibble -> odd index  */
            }
            if (n & 1) wf[n - 1] = K3_E2M1_PAIR[pb[(i0 / 2) + half]][0];

            double part = 0.0;
            for (int k = 0; k < n; k++) part += (double)x[i0 + k] * (double)wf[k];
            acc += part * (double)s;
        }
        y[o] = (float)acc;
    }
}

有两个细节值得停下来思考。if (s == 0.0f) continue 处理缩放字节 255，MXFP4 规范将其定义为 NaN；将其映射为零意味着一个损坏的字节会杀死一组 32 个权重，而不是将整行变成 NaN 并毒害所有下游层。而 if (n & 1) 处理元素数为奇数的组。对于 32 的组大小，这在当前检查点上永远不会发生，但还是写上了。这就是能工作的代码和正确的代码之间的区别，它只花了一行。

  PASS  mxfp4          64 rows x 3584 elems, EXACT on released checkpoint bytes

不是“在容差范围内”，而是精确。双方从相同的权重读取相同的字节，所以没有什么是可以合法不同的，测试以零容差编写以表明这一点。

这就是缩减一。专家以每个权重 0.53125 字节而不是 2 字节提供，将 5.45 TB 的专家权重降至 1.447 TB，并且永不扩展它们每 token 又节省了 194 GB 的内存流量。

5. 具有浮点契约的内核

每条代码路径必须产生相同的位。RMSNorm 是模型中使用最多的内核。

RMSNorm，epsilon 在平方根内，以 double 累加

void k3_rmsnorm(float *out, const float *x, const float *w, int n, float eps)
{
    double ss = 0.0;                                   /* double, not float */
    for (int i = 0; i < n; i++) ss += (double)x[i] * (double)x[i];
    const float inv = (float)(1.0 / sqrt(ss / (double)n + (double)eps));
    for (int i = 0; i < n; i++) out[i] = x[i] * inv * (w ? w[i] : 1.0f);
}

两个细节是承重的：累加器是 double，尽管每个输入和输出都是 float，并且 epsilon 放在平方根内部而不是外部。

固定的归约顺序，所以标量和 AVX2 逐位一致

/* Four accumulators partitioned by i % 4. This pins the summation order, so a
 * 4-wide vector loop adds the same numbers in the same sequence. */
static float k3_dot4(const float *x, const float *w, int n)
{
    double a0 = 0.0, a1 = 0.0, a2 = 0.0, a3 = 0.0;
    int i = 0;
    for (; i + 3 < n; i += 4) {
        a0 += (double)x[i]     * (double)w[i];
        a1 += (double)x[i + 1] * (double)w[i + 1];
        a2 += (double)x[i + 2] * (double)w[i + 2];
        a3 += (double)x[i + 3] * (double)w[i + 3];
    }
    for (; i < n; i++) a0 += (double)x[i] * (double)w[i];
    return (float)((a0 + a1) + (a2 + a3));       /* the parentheses are the contract */
}

bf16 到 fp32 是移位，不是转换，所以扩展是无损的

bfloat16 值是 float32 的高 16 位，低 16 位被丢弃，所以扩展是左移 16 位，并且是精确的。这就是为什么主干可以以其发布的精度流式传输，而没有任何精度问题，这一点在很久以后变得重要。

void k3_matmul_bf16(float *y, const float *x, const uint16_t *W, int in, int out)
{
#pragma omp parallel for schedule(static) if (out > 64)
    for (int o = 0; o < out; o++) {
        const uint16_t *row = W + (size_t)o * in;
        int i = 0;
        double acc;
#if defined(__AVX2__)
        {
            __m256d v = _mm256_setzero_pd();
            for (; i + 3 < in; i += 4) {
                /* bf16 -> f32 is a 16-bit shift: no table, no rounding */
                const __m128i h   = _mm_loadl_epi64((const __m128i *)(row + i));
                const __m128i b32 = _mm_slli_epi32(_mm_cvtepu16_epi32(h), 16);
                const __m256d wd  = _mm256_cvtps_pd(_mm_castsi128_ps(b32));
                const __m256d xd  = _mm256_cvtps_pd(_mm_loadu_ps(x + i));
                v = _mm256_add_pd(v, _mm256_mul_pd(wd, xd));   /* NOT fmadd */
            }
            double a[4];
            _mm256_storeu_pd(a, v);
            acc = (a[0] + a[1]) + (a[2] + a[3]);
        }
#else
        {
            double a0 = 0.0, a1 = 0.0, a2 = 0.0, a3 = 0.0;
            for (; i + 3 < in; i += 4) {
                a0 += (double)k3_bf16f(row[i    ]) * (double)x[i    ];
                a1 += (double)k3_bf16f(row[i + 1]) * (double)x[i + 1];
                a2 += (double)k3_bf16f(row[i + 2]) * (double)x[i + 2];
                a3 += (double)k3_bf16f(row[i + 3]) * (double)x[i + 3];
            }
            acc = (a0 + a1) + (a2 + a3);
        }
#endif
        for (; i < in; i++) acc += (double)k3_bf16f(row[i]) * (double)x[i];
        y[o] = (float)acc;
    }
}

两个分支都累加到四个 double 中，都按 i % 4 分区，并且都归约为 (a0 + a1) + (a2 + a3)。向量路径是标量路径，以相同顺序执行相同的加法，一次四个。

注意乘加。_mm256_add_pd 的 _mm256_mul_pd 故意不是 _mm256_fmadd_pd。融合乘加只舍入一次而不是两次，因此更精确，这正是问题所在：它会给出与标量循环不同的答案，硬件能力绝不能改变输出。

相同的权重，三条代码路径，一个哈希

证明 AVX2 和标量路径一致不能通过容差检查来完成，因为容差检查会愉快地通过一个悄悄重新关联其总和的核