第 21 章 改造手册:M0–M4
MoE 大模型低成本部署框架第 21 章 改造手册:M0–M4
CHAPTER 21

第 21 章 改造手册:M0–M4

11 分钟阅读 · 2,978 字

20 章把病灶定位到了四行代码和三个常量。这一章是动刀的顺序:
五个里程碑(M0 零代码 → M4 行预算扩容),每一步都能单独验收、单独回滚。

本书写作时(2026-10-21),M0–M4 一步都没有实施。 这一章给的是路径与判据,不是战果。
唯一例外是 §21.2 的 ①(网页默认贪心)——它已获决策批准,但尚未上线。


21.1 总原则与红线

四条,按被违反的代价排序:

  1. 服务对外在线。 任何停服操作(改引擎配置、换二进制 = 重启)必须先经用户同意。
    这是红线 1(17 章)。重启一条 lane 是 6–13 分钟(16 章),期间该 lane 的流量要靠网关 failover 接走。
  2. 只在 lane A(:8091,快道)上做实验。 lane B 是普道(:8092),压测只打 8091。
    动 lane B = 动公网流量的一半。
  3. 一次只改一个变量。 改前留 30 分钟基线,改后同负载对比。负结果也必须记下来
    ——负结果告诉你"瓶颈不在这一层",这比正结果更能省后面的时间。
  4. 不新写测量工具。 现成的四件套(20.8、13 章)已经够:
    step-stats-*.jsonl(逐窗口)、strata batch: 行(分项)、decode expert cache hit rate(命中率)、
    ~/bench(客户端口径)。

21.2 M0:零代码开关(只改配置 / 环境变量 / 前端默认值)

判据:4 槽 chain 每请求 ≥ 100 tok/s(2026-10-21 基线 67)。

① 网页默认贪心(决策已批准,未上线)

  • 依据:20.7 ②。temperature 0.6 ⇒ sl.sampled ⇒ 整组 return one_token()。
    实测 chain 194 vs batch 82 tok/s(1 槽)。
  • 做法:前端默认 temperature: 0(或只在"对话"预设里设 0),保留用户可调。
  • 验收:step-stats-a.jsonl 里 kind=chain 占比从 ~37% 升到 >80%;rows/slots 从 1.00 升到 ≥2。
  • 代价:0。不重启引擎,不改 C++。这是全书性价比最高的一处改动。

② --pcie-frac 1.0(排在 M1-2 之后才有意义)

  • 依据:PCIe 档每专家 ~70 µs 且不占 CPU;CPU 档每对 ~95 µs 且占核。现在只有 55% 的未命中走快档。
  • 风险:PCIe 流量从 ~486 MB/窗口涨到 ~970 MB/窗口(≈37 ms @ 22.7 GB/s)——可能不升反降。
    而且真正卡住的是 kStagingBlobs = 16:不改这个,pcie-frac 1.0 送不过去。
  • 验收:per layer-window 的 CPU experts 从 5.38 降到 ≤2;窗口时间下降。

③ STRATA_PF_FUSED=1

  • 依据:本包的量化类型全层都被融合 kernel 覆盖(moe_fused_iq.cu:1089 的 gu 集合、:1095 的 d 集合;
    4090D 的 cc=8.9 满足 moe_fused.cu:415),只差 moe_fused.cu:428 那个显式 opt-in
    (requested() 读 STRATA_PF_FUSED)。生产 profile 没设,日志里也确实没有
    prompt experts on the fused int8 kernels 这一行。
  • 预期:prefill 去掉每层一次 cudaStreamSynchronize(prefill.cpp:2055)。
    上游注释在原生包上记的收益不大(+8%/0%),但它把"prefill 每层卡主机"这件事消掉了,是 M3 的前提。
  • 风险:数值路径不同。必须核对同 prompt 贪心输出一致(前 32 个 token 逐个比)。

④ --short-read 512

  • 依据:read_windows(generate.cpp:6554)已经实现"用解码窗口读 prompt",
    只是被 --short-read 64(generate.cpp:459,门在 :6547,调用在 :6826)限死在 64 token。
  • 做法:把上限放到 256/512,让更多段走窗口路径。
  • 验收:kind=chain 窗口的行数上升,TTFT 不恶化。
  • 意义:这是唯一一个不改代码就能验证 M3 方向收益的实验。

已否决:--pool-workers 13 → 24

这条必须留在这里,因为它看起来太合理了,下一个人一定会再想到它。

表面推理 76% 的时间在 CPU 专家池,配额 32 CPU,而 worker 只有 13 → 加到 24
真实约束 run-iq3_xxs-a.sh 里 TASKSET=0-15(B 是 38-53),taskset -c 0-15 把 lane 钉死 16 核
算术 13 worker + 1 宿主线程 = 14 线程 / 16 核,已接近饱和;两条 lane 合计正好等于 cgroup 的 32 CPU
历史实证 2026-10-07 做过 24 worker,两个引擎 worker 在共享核上互踩,solo 掉到 9–17 tok/s(18.2)

⇒ CPU 专家这段只能"少做"(提命中率)或"换档"(走 PCIe/GPU),不能靠加线程。
这条否决直接把 M1 从"第二步"提成了第一优先。

这次的教训比结论更值钱:cgroup 配额不等于进程可用核数。
在容器里谈"加线程",先 grep TASKSET 看有没有人把进程钉住了。


21.3 M1:命中率与快档扩容(小改,只动常量与分配)

判据:命中率基线 +5pp;strata batch: 的 CPU experts 从 5.38 降到 ≤3.5。

M1-1 让显存专家缓存更满

现状 --expert-cache auto 算出 6434 槽 / 10.42 GiB,命中 63–92%(随负载波动;日志当日 88.6%)。
每 +1pp 命中 ≈ −0.3 ms/token ≈ +3% 吞吐(20 章 20.6 的斜率来源)。
手段:--vram-reserve-mib 1600 略降、别让别的缓冲占住显存。收益确定,但天花板受显存限制。

M1-2 kStagingBlobs 16 → 32(第一处要改的代码)

// include/strata/core/verify.hpp:357
static constexpr int64_t kStagingBlobs = 16;   // 暂存区可容纳的专家块数
  • 成本:32 × 1.87 MB ≈ +30 MB 显存(启动剩 1182 MiB)。改到 64 是 +90 MB。
  • 为什么只加 30 MB:暂存区在 Verifier::init(verify.cpp:336)里按 kStagingBlobs × max_blob
    只分配一次,每个 Verifier 一份、不是每个 batch 会话一份。
  • 风险:① arena_ 总量(count.used)要与 --expert-cache auto 的余量复核;
    ② 多卡(--layer-split)时 staging_cap 会按 stage 再切(verify.cpp:864/1326 的 per = G==2 ? cap/2 : cap),
    单卡不受影响。
  • 验收:per layer-window 的 PCIe 计数上升、CPU experts 下降;窗口 ms 下降。
  • 回滚:改回 16 重编;二进制走 switch_engine.py dev build 的备份,不覆盖生产。

M1-3 让"下一 token 的专家"提前进显存

RouterLookahead(expert_source.cpp:1904-2038)已在代码里,但只在 src.warms() 时启用
(generate.cpp:3673-3684),即只对文件层原地映射有效;本项目用的是锁页 arena 包,
大概率没开(待核实:重启后日志里找 routing-aware prefetch)。
若没开,这条路是"改一行条件"级别的改动,值得先确认再谈实现。


21.4 M2:拆掉每层主机屏障(核心改动)

判据:窗口里 CUDA0 GPU-reach wait 占比从 23% 降到 <10%。

三个层次,风险递增,做的顺序也是这个顺序:

层 改法 风险 收益上限
M2-1 下一层预取 主机算第 l 层 CPU 专家时,用路由器权重先算好第 l+1 层要哪些专家并预取 低(不动 GPU 侧协议) 把 70 ms 藏进 21 ms(上限),实际看能盖住多少
M2-2 两窗口流水 两个 CUDA Graph 实例交替,用两条 stream 让 GPU 在上一窗口还有活时领下一窗口 中(显存 ×2、*seq/*flag 要按窗口编号区分) 门铃等待被彻底隐藏
M2-3 CPU 专家搬到 GPU 未命中专家本来就在锁页主机内存里,用 cudaHostAllocMapped 零拷贝让 GPU 直接读 中高(延迟高但不占 CPU) 与 M2-1 正交,可叠加

M2-1 的前提已经具备:路由器是 ffn_gate_inp.weight(BF16,n_expert × n_embd × 2 = 2.6 MB/层),
48 层全部常驻主机只要 125 MB。


21.5 M3:decode 与 prefill 同窗(最初那个设想)

判据:解码槽最长停顿 ≤0.5 s(现 4.9 s);prefill 总吞吐不掉 >5%。

这一章要写清楚一个容易被误解的点:"逐层交错"不是靠给算子加锁得到的。

  • 同一请求的层间是强依赖(第 l+1 层要第 l 层的输出),单请求内部没有可并行空间;
  • 可并行的是不同请求,而它们现在已经共享窗口——卡住的是窗口粒度是整窗、且行预算被并发除(20.7 ①);
  • prefill 要能"插进"解码,前提是 prefill 也走窗口,否则它是一个 2.4 s 的同步调用,中间插不进任何东西。

三步:

步 做什么 依据
M3-1 把 sp.run()(同步,generate.cpp:6831/6837)拆成"每轮推一个 chunk"的状态机 等价于把 STRATA_PREFILL_YIELD_CHUNKS 从 2 降到 0.5,但不再需要 refill/lend
M3-2 把 prompt 行编成一个"伪槽",与真实槽一起进窗口 行里已有独立 pos/token,数据结构不用改;要改的是"每槽一行/多行"的语义 → “每槽一个行数向量”
M3-3 打开 --short-read 的适用范围 见 21.2 ④,零代码先验证

M3-2 真正的坑在 kernel 上限:每专家组最多 8 条目(GMAX),prefill 行 + decode 行会超(M4)。

收益的正确预期:不增加 GPU 总工作量,所以吞吐不变、延迟分布变好。
这与 Sarathi-Serve 的结论一致(docs/references/papers/sarathi-serve-osdi24.pdf,本机已存全书 12 篇论文)。


21.6 M4:行预算扩容(8 → 16/32)——最后做

判据:8 槽并发时每请求 ≥ 60% 单路速度。

为什么最后做:它要同时改三处 kernel + 重录所有 CUDA Graph,而收益依赖前面的命中率/CPU 项先解决。

要动的六处:

  1. kVerifyMaxT verify_kernels.hpp:28(8 → 16)
  2. MAXT kq_avx2.cpp:43 + expert.hpp:169
  3. GMAX s2_expert_grouped.cu:1360——并去掉 min(..., GMAX) 的静默截断(改成显式报错或分多组)
  4. --batch 上限 generate.cpp:2960-2961
  5. dma_src[64](expert_source.cpp:3384)与 staging_cap 的联动
  6. 所有 capture_batch(base, S, L) 的键空间(S、L 组合变多 ⇒ 录制时间与显存上涨)

收益模型:L = min(4, kVerifyMaxT/S),所以扩到 16 之后 S=4 仍是 4 行
(chain_rows_per_slot = 4 才是真上限)。扩 kVerifyMaxT 只对 S≥3 有用,且要先放开 chain_rows_per_slot。

更划算的替代:保持窗口 8 行,但让每个槽独立决定这一轮参不参与(M2-2 的副产品),
避免"3 个槽在跑、第 4 个闲着也占份额"的浪费。先做这个,M4 可能就不需要了。


21.7 验收口径与实验协议

(沿用 20.8;表里每行都能用现成留档算出来,不需要新工具。)

指标 来源 备注
每窗口时间 / 行数 / 槽数 ~/logs/step-stats-a.jsonl kind in (batch, chain)
投机是否生效 同上 drafts / accepted 判 20.7 ② 的唯一直接证据
prefill 吞吐 同上 kind=prefill(generate.cpp:6600 出口)
命中率 decode expert cache hit rate: xx.x% 引擎日志
窗口分项 strata batch: 行 槽全空时打印一次
每请求 tok/s 上面组合 每槽行数 ÷ 窗口 ms × 1000
端到端 ~/bench/ 现有压测台 只打 8091

协议:① 一次一个变量,改前 30 分钟基线;② 重启 lane A 前先问;
③ 压测前看 /shared/pressure/pressure.json,任一 high 就退避;
④ 每个实验一条 ~/bin/note;⑤ 负结果也记。


21.8 改代码的工作流(本项目实况)

引擎代码不直接改生产二进制。现成的工作流是 SSEStraBed/deploy/switch_engine.py:

python3 deploy/switch_engine.py status                     # 两条 lane 现在指哪个二进制
python3 deploy/switch_engine.py dev build --lane a         # 构建到该 lane 的 build/,写 BUILD.json
python3 deploy/switch_engine.py dev run   --lane a --yes   # 切过去(= 停服 + 拉起,需授权)
python3 deploy/switch_engine.py dev rollback --lane a      # 切回上一个

三条必须知道的:

  1. 每个 lane 一套 build/,版本号写在 BUILD.json(engine_build)。换二进制前先看 status。
  2. dev run 就是一次停服(16 章)。默认 lane 是 b——不写 --lane a 会打到普道。
  3. 420 秒内起不来的引擎会被 supervisord 判死(16 章的时间账),所以改常量这种小改动也要先确认能编过再切。

21.9 本章小结

  • 顺序:贪心默认(0 成本)→ kStagingBlobs 32(+30 MB,直击 76%)→ pcie-frac 配合 →
    STRATA_PF_FUSED=1(prefill 去每层同步)→ --short-read 512(M3 原型)→ M2 → M3 → M4。
  • M4 可能不需要:先让每个槽独立决定参不参与,比扩 kernel 便宜一个数量级。
  • --pool-workers 不要动:taskset 已经把它钉在 16 核上,13 是饱和值。
  • 每一步都要能单独回滚;每个实验写一条 note,负结果同样记。
  • 这本书写作时一步都没实施——19.1 那张表里 #1/#3/#9 仍然是"未修"。

本章出处

内容 出处
M0–M4 的定义与判据 docs/plans/ENGINE-CODEBASE-AND-OPTIMIZATION-PLAN-2026-10-21.md §15–§20
kStagingBlobs / 门铃 / L_max / sampled 20 章出处表(同一批源码行)
STRATA_PF_FUSED opt-in src/prefill/moe_fused.cu:415/421/428;moe_fused_iq.cu:1089/1095
read_windows 与 --short-read generate.cpp:459/6547/6554/6826
taskset 钉核与 24 worker 事故 run-iq3_xxs-{a,b}.sh;18.2
switch_engine.py 子命令 SSEStraBed/deploy/switch_engine.py(`dev build
变长窗口实验的结论 10.3.1;docs/reports/DEEPINFER-VARLEN-2026-10-21.md
全部章节