第 21 章 改造手册:M0–M4
20 章把病灶定位到了四行代码和三个常量。这一章是动刀的顺序:
五个里程碑(M0 零代码 → M4 行预算扩容),每一步都能单独验收、单独回滚。本书写作时(2026-10-21),M0–M4 一步都没有实施。 这一章给的是路径与判据,不是战果。
唯一例外是 §21.2 的 ①(网页默认贪心)——它已获决策批准,但尚未上线。
21.1 总原则与红线
四条,按被违反的代价排序:
- 服务对外在线。 任何停服操作(改引擎配置、换二进制 = 重启)必须先经用户同意。
这是红线 1(17 章)。重启一条 lane 是 6–13 分钟(16 章),期间该 lane 的流量要靠网关 failover 接走。 - 只在 lane A(:8091,快道)上做实验。 lane B 是普道(:8092),压测只打 8091。
动 lane B = 动公网流量的一半。 - 一次只改一个变量。 改前留 30 分钟基线,改后同负载对比。负结果也必须记下来
——负结果告诉你"瓶颈不在这一层",这比正结果更能省后面的时间。 - 不新写测量工具。 现成的四件套(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
--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 项先解决。
要动的六处:
kVerifyMaxTverify_kernels.hpp:28(8 → 16)MAXTkq_avx2.cpp:43+expert.hpp:169GMAXs2_expert_grouped.cu:1360——并去掉min(..., GMAX)的静默截断(改成显式报错或分多组)--batch上限generate.cpp:2960-2961dma_src[64](expert_source.cpp:3384)与staging_cap的联动- 所有
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 # 切回上一个
三条必须知道的:
- 每个 lane 一套
build/,版本号写在BUILD.json(engine_build)。换二进制前先看status。 dev run就是一次停服(16 章)。默认 lane 是 b——不写--lane a会打到普道。- 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 |