第 10 章 方法学：怎么确认主矛盾真的在说的地方
一套异构系统里什么都能是瓶颈。唯一靠谱的方式是把每一类可能性逐一排掉——不靠感觉，靠测量。
10.1 问题意识：为什么不能靠感觉
在 duck-llm 的开发过程里，性能判断从来不是只靠“跑起来感觉变快了”来下结论的。一个异构系统的延迟由多层因素叠加：GPU 算力、PCIe 传输、网络序列化、duck 侧内存带宽、duck 侧指令流水、Python GIL 等等。每一层都可以成为限制因素，而且这些因素之间经常互相遮掩——压掉一个，另一个才会浮出来。
这章要说的，是这套系统在关键决策点上用了哪些方法来确认“主矛盾在哪里”。
具体来说，方法学最密集的地方是 2026 年 3 月下旬 fp9 kernel 调优阶段。那时系统已经跑通 DeepSeek-V3.2，但不清楚 MoE forward 的吞吐瓶颈在哪里——是内存带宽打满了，还是指令成本还高？这个判断直接决定下一步该投入多少精力在 kernel 微调上。
10.2 字节口径：fp9 每个参数读多少字节
先把单位算清楚
在做任何 roofline 估算之前，必须先把“fp9 每个参数到底读多少字节”算清楚。这不是废话——fp9 格式同时有三段 payload，搞混了比例就会让带宽利用率的估算差几成。
当前 exact64 down_proj 每个参数的 payload 布局是：
weight_main：1 B/param（主权重，9-bit 存到 1 字节，高位即 sign 位）sign：1/8 B/param（额外符号位，每 8 个参数共享 1 字节）scale：1/2048 B/param（每 128 参数一个 fp32 scale，后来 64/32 各自有更细的 scale block）
合起来：约 1.126 B/param，即约 9.004 bit/param。
这个换算系数是后续所有 GB/s → Gparam/s 折算的基准。
真实 MoE hot 的吞吐
2026-03-30 的 breakdown 实测，exact64_combo moeish-hot, selected_experts=20, 4-core 的 wall 约折算为 20.30 GB/s，kernel-max（不含 finalize 等后处理）约折算为 22.05 GB/s。
这两个数字代表的是：在真实 MoE forward 里，fp9 kernel 实际消耗内存带宽的量级。
10.3 roofline 测量：duck 上的 effective read ceiling 是多少
下文多处出现“roofline”一词。它是计算机体系结构性能分析中的常用术语，含义是“给定硬件条件下某类工作负载的理论吞吐上限”。本文借用这个思路，不是画经典的 roofline 图，而是直接在目标平台上测出 read-only 的有效上限。
为什么需要专门测
“32GB DDR4-3200 理论带宽 25.6 GB/s”这个数字是理论值，不能直接当作 roofline。实际系统里还有 L1/L2/L3 cache 层次、prefetch、store buffer 等因素，真实 read-only 的有效上限未必等于这个理论值，也可能因为平台差异而不同。
因此，需要在这台 N100 上直接测。
harness 设计
2026-03-30 新建了两份专门的 roofline harness：
- 单核版本：
bench_memread_roofline_single.cpp - 4 核 same-ELF 版本：
bench_memread_roofline_smp.cpp
两份 harness 内部都带自己的 sweep 逻辑，不在外部并行触发（避免两个 bench 抢同一块内存路径）。
read loop 的实现用了 AVX2 指令：
_mm256_load_si256(...)
_mm256_xor_si256(...)并通过把结果落到 volatile 全局变量，避免编译器把整段 loop 当成无副作用代码优化掉。
working set 设计了两种模式：
- linear：连续大块读，working set 从 32MB 扫到 384MB
- payload：按真实 MoE expert 的 layout 分组读（
weight_main + sign + scale，单个 expert 约 516 KB），层数从 1 到 24 sweep
每个 sweep 点都选不同 offset，避免重复读同一段带来的 cache residue 问题。
结果
| 模式 | 单核 p50 | 4 核 p50 |
|---|---|---|
| linear 384MB | ~17.00 GB/s | ~35.67 GB/s |
| payload 24 layers | ~16.97 GB/s | ~34.93 GB/s |
从 64MB working set 开始，结果就已经非常稳定。linear 和 payload 两种模式收敛到几乎同一档，说明 MoE 那种分散读法和连续 linear 读在这个规模下没有明显区别。
4 核相对单核的倍数：约 2.06~2.10x，不是理想的 4x。
10.4 方法学复核：这些数字可信吗
这组结果比较反直觉——4 核 ~35 GB/s 的 aggregate 结果明显高于 DDR4-3200 单通道理论的 25.6 GB/s。有四类可能的解释需要逐一排除：
一、4 核有没有重复读同一区域
直接查了 harness 代码里的地址切分逻辑：
- linear：4 个核按
cpu_idx均分总 working set，每核读 disjoint slice（例如 384MB 总量，每核 96MB） - payload：按 layer rank 切分，不同核处理不同 layer，没有跨核重复
因此 4 核 roofline 不是靠“重复读同一区域抬出来的”。
二、GCC 有没有把 read loop 优化掉
直接用 objdump 看了 bench_memread_roofline_single.elf 的反汇编：
vpxor (%rax), %ymm0, %ymm0
vmovdqa ...可以看到对 g_weight_main_64、g_weight_sign_64、g_scale_bytes_64 的真实内存 load 指令，以及 payload scan 路上的 loop 结构。read loop 没有被“优化成纯算术或常量折叠”。
三、17 GB/s 这个量级是否合理
专门补了 libc sanity bench，用最小程序分别跑：
- 单核
memset(384MB)：~9.86 GB/s - 单核
memcpy(384MB)：~7.09 GB/s - 4 核
memset(384MB)：~29.75 GB/s - 4 核
memcpy(384MB)：~20.58 GB/s
这组数据说明了一件事：read-only 明显比 memset/memcpy 高，是正常现象。memset 是纯写，经过 write-allocate / store buffer；memcpy 同时涉及读和写；而 read-only 只有读流量，路径最短。所以“read-only > memset > memcpy”这个排序本来就合理，17 GB/s 并不像“明显测错了”。
四、forced-cold：把窗口真正换成 unique
上面的实验用了 rotating window（窗口轮转），仍然可能有“后面的 trial 回访了旧窗口”的疑点。
2026-03-30 做了两轮强化版强制冷读：
- rotating non-overlap windows：backing store 拉大到 768MB（linear）和 473MB（payload），每个 trial 切到新的不重叠窗口
- true unique windows：把
measure_trials = min(requested_measure_trials, windows)，warmup 设 0，保证每个 sweep 点只测真正不重复的窗口
两轮结果几乎完全不变：
| 模式 | 单核 p50 | 4 核 p50 |
|---|---|---|
| forced-cold linear 64MB | ~17.02 GB/s | ~34.94~35.56 GB/s |
| unique-window payload 4 layers | ~17.12 GB/s | ~35.03~35.06 GB/s |
“同一窗口 cache 残留”和“窗口回访”这两条解释，到这一步都被排掉了。
剩下的 caveat
排掉上述四类可能的测量误差之后，一个矛盾仍然没有消失：这组 effective read ceiling（单核 ~17 GB/s，4 核 ~35 GB/s）比 N100 硬件前提（DDR4-3200 单通道 ~25.6 GB/s）更高。
因此，这套数字最准确的口径是：平台上 effective logical read ceiling，而不是 raw DDR pin bandwidth。后者需要更底层的测量手段（如 IMC 计数器、STREAM benchmark），目前还没有做。这个 caveat 在报告里留着。
10.5 主矛盾定位：MoE fp9 还远没打满带宽上限
有了 effective read ceiling，现在能做的定量判断是：
- MoE fp9 hot (wall)：约 20.30 GB/s → 4 核 effective ceiling ~35 GB/s 的 58%
- MoE fp9 hot (kernel-max)：约 22.05 GB/s → 4 核 effective ceiling 的 63%
这说明“内存带宽天花板已经打满”不是当前的瓶颈解释。真正拖住 Gparam/s 的，是：
- fp9 decode/dot 本体的每字节指令成本（AVX2 整数/浮点混合路径）
slot-scale：每层 MoE 完成后，按 topk expert 权重对各核贡献做缩放合并finalize-accum：CPU0 收各核结果做 fp32 累加
在实测 breakdown 里（selected_experts=20, 4-core hot），kernel-max 约 468 ms，slot-scale-max 约 16.8 ms，finalize-accum 约 21.6 ms。后两项合占 wall 约 9%，是值得继续压缩的代价。
10.6 专家数 sweep：为什么 Gparam/s 随专家数增加而提升
一个反直觉现象
做 MoE forward 时，如果 selected_experts 从 4 增加到 20（参数量 5x），按直觉应该 wall time 也是 5x。但实测并不是这样：
| selected_experts | cold wall p50 (ns) | cold Gparam/s |
|---|---|---|
| 4 | 131,484 | 13.96 |
| 8 | 231,560 | 15.85 |
| 12 | 329,988 | 16.68 |
| 16 | 428,790 | 17.12 |
| 20 | 530,443 | 17.30 |
参数量 5x，wall time 只涨 4x，Gparam/s 反而更高。
原因
这里有一个固定成本的摊销效应：不管 selected_experts 是 4 还是 20，barrier 同步、worker 起跑时延、start-max-delay、worker-tail-gap 这些固定控制面的开销基本不变（各自约几百 ns 到几微秒）。expert 数越多，kernel 本体的运行时间越长，这些固定成本被摊薄，Gparam/s 自然往上爬。
真实 DeepSeek-V3.2 MoE 每 token 路由 8 个 expert（top_k=8），对应更接近 selected_experts=8 这一档，hot Gparam/s 约 16.6。如果每步有机会同时处理更多 expert（如 MTP 情况下 bsz=2），摊销效果更明显。
10.7 prefetch 实验：为什么整包 prefetch expert payload 是负收益
直觉和预期
一个自然的优化想法是：在每个 expert forward 之前，先用 software prefetch 把这个 expert 的全部 payload（weight_main + sign + scale，约 516 KB）预取进 L3 cache，再开始计算。
测了一下。
结果
在 bench_fp9_moeish_exact64_smp.cpp 里新增 prefetch_mode=1，对每个 expert 的 payload 按 64 字节 stride 全量做 _mm_prefetch(..., _MM_HINT_T0)。结果：
| selected_experts | baseline hot wall | prefetch hot wall | 变化 |
|---|---|---|---|
| 8 | ~221ms | 更慢 | 负收益 |
| 20 | ~511ms | 更慢 | 负收益，且 experts 越多越差 |
prefetch 路本质上是“额外再扫一遍 ~516 KB/expert 的 payload”。这条扫描本身消耗带宽，也可能污染 L3 cache（把其他正在用的数据挤出去），而真正的 forward 计算本身已经能驱动硬件 prefetch 跟上，反而不需要额外的 software prefetch 帮忙。
这个实验排掉了一条看起来合理的优化方向，避免在后续 kernel 工作里投入无效精力。
10.8 trace 与 breakdown：作为判断支撑的作用
端到端 breakdown
在 decode 路径调优时，breakdown 信息是确认“固定成本在哪里”的主要工具。2026-03-12 对 Qwen3-30B-A3B decode 做的 breakdown 揭示：约 29ms/token 的非 duck 时间里，attention 的 torch.cat KV cache 追加是最大的单项。这个判断直接驱动了 static KV cache + flash attention 2 的方案选择，最终带来 flattened runtime 的整体改造。
MoE forward breakdown
moeish-hot 的 trial-level breakdown 把 wall 拆成了：
start-max-delay：BSP 发起到所有 worker 就位的最大 skewworker-tail-gap：最后一个 worker 完成到 barrier unlock 的延迟kernel-max：4 核中 kernel 本体耗时最长的那一核slot-scale-max：最慢核完成 slot 缩放的时间finalize-zero / finalize-accum：CPU0 做最终归并的时间
这种精细 breakdown 的作用是：避免把“整体 wall 时间”归因于某一类成本。例如，start-max-delay 和 worker-tail-gap 加起来其实很小（~几百 ns 级），不是主要成本；真正占重的是 kernel-max 和 finalize-accum。
MTP1 dual-lane 的 overlap 摘要
第 9 章的 pair overlap runtime 实现了 per-step 的 overlap 摘要日志：
olap pair=112ms serial=199ms gain=87ms ratio=0.43 s0=...ms s1=...ms wait=.../...ms这条信息在 549 个连续 step 的长跑里持续打印，确认了 dual-lane overlap 不是偶发现象，而是稳定在 ~43% 的 gain ratio。没有这条实时 breakdown，很难有把握说 overlap 在任意 prompt 上都是正收益。
10.9 这套方法学的核心价值
把这章的内容总结成一句话：这套优化不是靠“感觉变快了”来下结论的。
每一个关键决策背后都有可以反驳的数字：roofline 测了，复核做了，libc sanity 对过了，forced-cold 再跑了。prefetch 看起来合理，测了发现是负收益。整包优化看起来没空间，但 breakdown 告诉你固定成本还有多少。
这种思路的代价是：要写更多 harness，做更多 ablation，而不是直接把优化往 active 代码里塞。但它带来的好处也很实在：不会被“这次跑快了但原因不明”的结果骗走时间，也更容易向别人（或者半年后的自己）解释为什么结论是成立的。
项目声明 / Project Disclaimer
本项目为作者以个人身份、利用业余时间推进的个人娱乐项目；除非另有明确说明，它与作者的任何雇主、客户、学校、单位或其他组织均无合作、雇佣、委托、赞助或背书关系，也不代表任何该等主体的立场。除普通个人捐赠外，本项目未获得任何资金支持。
This project is a personal hobby project developed by the author in a personal capacity and in personal time. Unless explicitly stated otherwise, it has no collaboration, employment, commission, sponsorship, endorsement, or institutional affiliation with any employer, client, school, partner organization, or other entity, and it does not represent any such party's views. No funding was received for this project except ordinary personal donations.
许可 / License
除非另有说明，本页原创文字、本站原创图片与本站原创图表采用 CC BY-NC-ND 4.0 发布。
转载时请保留原文标题、署名“JudgeDuck AI”、发布日期与原始链接；禁止商业转载、改写、摘编、翻译或基于原文创作演绎作品。
第三方商标、外部链接内容，以及文中另有标注的材料，不在上述许可范围内。
Unless otherwise noted, the original text, original images, and original figures on this page are licensed under the Creative Commons Attribution-NonCommercial-NoDerivatives 4.0 International License.