第 20 章 解码窗口的解剖：那 92.5 毫秒去哪了
这一章是全书唯一一处把一帧解码拆到毫秒的地方，也是阅读 07/08/10 三章之后必须补上的一页。
它回答一个决定了后续所有改造方向的问题：并发变慢，到底是 GPU 不够快，还是别的什么不够快？结论先写在这里：在解码时，GPU 大部分时间是闲的。等它的是主机。
所以「提高 GPU 利用率」这个常见的优化直觉，在这套系统上从一开始就问错了问题。
20.1 一个反直觉的观察
先摆三个数字（lane A，2026-10-21）：
| 观察 | 数字 | 说明 |
|---|---|---|
| 单路解码 | 100–120 tok/s | 与上下文长度几乎无关（9.1） |
| 4 路并发聚合 | ~86–153 tok/s | 只比单路好一点，远不是 4 倍（19.1 #9） |
| 4 路时每路 | ~22–39 tok/s | 腰斩再腰斩 |
直觉的解释是"GPU 算不过来了"。但引擎自己打印的账否掉了这个解释。
一个重载样本（历史日志，__batch__ 汇总行）：
strata batch: ... avg 2.84 rows, 92.50 ms/window
= run 90.86 (CUDA0 GPU-reach wait 21.14 + CPU experts 70.00)
+ commit 1.63 + emit ...
per layer-window: CPU experts 5.38, VRAM hits 18.36, PCIe 5.41
把这一行翻译成人话：
92.5 ms 的一个解码窗口（约 3 行，即并发 3 路各出 1 个 token）
├─ 70.0 ms CPU 专家：主机上 13 个 AVX2 线程在算"显存里没有的专家" ← 76%
├─ 21.1 ms GPU 空转：卡在门铃上，等主机把手上的活干完 ← 23%
└─ 1.6 ms 提交/输出（commit + emit）
GPU 在这 92.5 ms 里只真正干了不到 1.6 ms 的"自己的活"——其余时间它在等人。
这就是为什么「换更快的卡」不会让这套系统变快（19.5 已否决清单里那条"换更小量化冲并发"的更深一层原因）。
20.2 窗口的一生：从 BGEN 到 BDONE
一次解码不是一个 token 一步，而是一帧窗口。它的代码路径（SSEStraBed，8177 行口径，2026-10-21 复核）：
| 步 | 位置 | 做什么 |
|---|---|---|
| 1 | generate.cpp:6075 |
主循环 for (;;) 醒来；全引擎只有一个线程在这里转 |
| 2 | generate.cpp:6076 |
有活跃槽就进 chain_step()（--batch-mtp）或 batch_step() |
| 3 | generate.cpp:5781 |
算这一帧的行预算 L_max（见 20.7 ①） |
| 4 | generate.cpp:5823 |
检查组里有没有"非贪心槽"（见 20.7 ②） |
| 5 | verify.cpp:1755 |
run_slots_chain(base, S, L, …)：S 个槽 × L 行，一次性录进一张 CUDA 图 |
| 6 | verify.cpp:411 |
record_window(T)：把这一帧所有 GPU 操作捕获成图（T = S × L） |
| 7 | verify.cpp:1768-1811 |
逐层门铃握手（20.3；非链路径同构，在 :1695-1738） |
| 8 | generate.cpp:5860（commit_chain）/ :5695（BDONE） |
按每槽验证结果提交 token，写 BT/BDONE |
关键在第 7 步：一帧之内，48 层是逐层流过 GPU 的，而每一层都要停下来跟主机对一次账。
20.3 门铃：每层一次主机↔设备握手
verify.cpp:1755-1827 的循环是全篇最该背下来的一段（中文注释是本地化补丁加的）：
const int64_t steps = le_ - lb_; // 本 stage 的层数（单卡 = 48）
for (int64_t k = 0; k < steps; ++k) {
const uint32_t want = (uint32_t)(k + 1);
while (*seq < want) { _mm_pause(); ... } // ① 主机自旋：等 GPU 爬到第 k 层
set_plan_slot(0);
if (pool != nullptr)
pool(user, h_x_, h_ids_, (int) R, ss_->k, h_ymiss_, l); // ② 同步跑完本层全部 R 行的 CPU 专家
std::atomic_thread_fence(std::memory_order_seq_cst);
*flag = want; // ③ 抬 flag，放行 GPU 进下一层
ms_wait += ...; ms_pool += ...; // ④ 分别记账（这俩就是上面那 21.14 / 70.00）
}
GPU 那一侧对应的是一个只占 1 个线程的等待核（src/kernels/cuda/verify_kernels.cu:504，
启动点 :616）：
wait_flag_ge_kernel<<<1, 1, 0, stream>>>(flag, value);
也就是说：GPU 跑到第 k 层，派一个线程去自旋等一个内存里的 flag，其余 SM 全空着；
主机把这一层的 miss 专家算完，写 flag；GPU 才继续。48 层 = 48 次这样的往返。
三点必须说清楚，否则容易误读：
- 这不是"GPU 慢"，是"GPU 在等一个必须由 CPU 先算完的东西"。 门控残差链是严格串行的
（07.1 的"96% 暴露"就是这个意思），第 l+1 层要用第 l 层的归属专家结果。 - 两次握手之间确实是重叠的：GPU 算第 k 层的 deep 部分时，主机可以已经在算第 k-1 层的专家。
所以wait与CPU experts不是总能直接相加（见 20.5 的第三个样本，51.21 + 31.99 > 41.85）。
但在 CPU 成为瓶颈的重载窗口里，它们接近相加（21.14 + 70.00 = 90.86 ≈ run 90.86）。 - 并发越高，② 里的
R越大（R = S × L），这一层的 CPU 活就越重，门铃等得越久。
这是"多一路就慢一截"最直接的机制。
20.4 怎么读那行 strata batch:（这里有个坑）
这行汇总由 generate.cpp:5711-5725 在所有槽都空下来时打印一次（batch_on() = 有任意槽活跃，
所以它出现在每段负载的尾巴上）。字段分两类，混读会得出完全错误的结论：
| 字段 | 单位 | 含义 |
|---|---|---|
avg N rows |
行 / token | 两条路径含义不同：batch_step 里是每窗口行数（bt_rows += S，generate.cpp:5683）；chain_step 里是每窗口提交的 token 数（每提交一个 ++bt_rows，:5894） |
X ms/window |
毫秒 | 一帧的总时长（= run + commit + emit） |
run / GPU-reach wait / CPU experts |
毫秒 | 主机侧记账（ms_wait / ms_pool） |
per layer-window: CPU experts / VRAM hits / PCIe |
计数，不是毫秒 | 每个（窗口 × 层）里各类专家各多少个 |
末尾 N rows/s |
行/秒 | 用上面的"行"算的，所以同一条链式汇总里它其实是 token/s |
第三、四行各有一个坑，可以凑成一条完整的教训：
- 第四行是计数，不是毫秒。 我第一次读时把
5.38 / 18.36 / 5.41当成毫秒，账怎么算都不对。 - 第一行在两条路径里是两个量。 于是"三者之和 ≈ 行数 × 10"这条校验只在
batch_step的行上成立
（重载样本：5.38+18.36+5.41 = 29.15 ≈ 2.84 × 10✓）；
在chain_step的行上不成立（下文 6.62 那个样本：7.33+132.32+6.47 = 146，
而真正的行数是S × L，最大 16，×10 = 160 ⇒ 其实是 ~91% 命中，
拿 6.62 去乘 10 会得出"命中 220%"这种荒谬结论）。
一般的教训：同一行输出里混着两种单位和两种语义时，作者有义务把它们写进格式串。
这行两者都没写，代价是每一个第一次读它的人都会错一次——包括写这本书的人。
20.5 三个样本，三种负载形态
同一行输出，在不同负载下形状差别极大。三个真实样本（都出自 lane A）：
| 样本 | 路径 | avg rows |
ms/窗口 | GPU-reach wait | CPU experts | commit | 形状 |
|---|---|---|---|---|---|---|---|
| 历史重载（cache 冷 / 并发 3） | batch（†） | 2.84 行 | 92.50 | 21.14 | 70.00 | 1.63 | CPU 主导，76/23 |
| 今日热态（并发 4，短上下文） | batch | 3.95 行 | 26.5–27.7 | 15.9 | 8.9–9.9 | 1.7 | 门铃与 CPU 各约 1/3 |
| 今日重窗口 | chain | 6.62 token | 41.9 | 51.21 | 31.99 | 0.03 | 等待 > 总时长 → 存在重叠 |
（†）第一行的路径是推出来的，不是日志里写的：它的 avg rows 与 per layer-window
能对上账（29.15 ≈ 2.84×10），而这条校验只在 batch_step 的行上成立（20.4）。
第三行值得单独看两件事：
15.9 + 8.9 ≈ 24.8 ≈ run 24.78（第一、二行几乎相加）；
但51.21 + 31.99 = 83.2 > 41.85。两个计时器互相重叠，既不能相加，也不能各自当瓶颈证据
——要看谁在挡住谁。- 它的
avg rows其实是提交 token 数（chain 路径），不能拿去和per layer-window对账（20.4）。
per layer-window 也说明命中率差别很大：热态样本 VRAM hits 36.30（几乎全命中），
重载样本只有 18.36，剩下的 5.38 + 5.41 个要走 CPU 与 PCIe。
命中率就是这章的主角（07 章讲过它的测量方法）。
20.6 三个硬上限：为什么并发一高就掉进 CPU 档
一个未命中的专家有三条去处（07 章的三级存储，代码里的决策在
src/core/expert_source.cpp:3376-3413，每层算一次）：
命中 VRAM → GPU 直接算 ← 最快
走 PCIe → 从锁页 arena DMA 到 staging 再算 ← 次快（也贵，但不动 CPU）
走 CPU → 主机 AVX2 池现算 ← 最慢，且在关键路径上
分的比例由 --pcie-frac 决定（默认基值 0.55，启动时按实测 PCIe 带宽再修正；
lane A 日志：PCIe probe: 22.7 GB/s host->device → pcie_frac 0.55）。
但真正卡住并发的是另外两个没人告诉你的上限：
| 上限 | 位置 | 后果 |
|---|---|---|
kStagingBlobs = 16 |
include/strata/core/verify.hpp:357 |
一帧最多 16 个 miss 专家能走 PCIe 快档，第 17 个起全部落回 CPU 档 |
dma_src[64] / fetches < 64 |
src/core/expert_source.cpp:3384 起 |
单层的 DMA 源指针数组硬上限 64 |
行数上限 kVerifyMaxT = 8 |
include/strata/kernels/verify_kernels.hpp:28；GMAX = 8 在 s2_expert_grouped.cu:1360 |
一帧总行数 > 8 会被静默截断（10.4 已讲） |
kStagingBlobs = 16 是这一章新加的发现，也是最"便宜"的一个瓶颈：
- 暂存区在
Verifier::init里只分配一次（verify.cpp:336的b.take<uint8_t>(kStagingBlobs * max_blob)，
verify.cpp:365是sink_.staging_cap = kStagingBlobs），单卡只有一份。 - 每个 blob ≈ 1.87 MB ⇒ 16 → 32 只是 +30 MB 显存（启动时还剩 1182 MiB）。
- 一帧 4 行 × 10 专家 = 40 个未命中，
--pcie-frac 0.55想送 22 个过去，实际只送出 16 个。
16这个数字不在任何文档里，只在源码里。
找它花的时间，比它将来省的时间少得多——这是"读源码而不是读文档"最具体的一次回报。
20.7 并发退化的四行代码
把上面所有机制压成四行可以指认的代码（不是"架构感觉"，是 grep -n 能定位的四行）：
| # | 代码 | 位置 | 后果 |
|---|---|---|---|
| ① | const int L_max = std::min(chain_rows_per_slot, std::max(1, (int)kVerifyMaxT / S)); |
generate.cpp:5781 |
4 路时投机额度 4→2；行预算被并发除掉 |
| ② | if (sl.active && sl.sampled) return one_token(); |
generate.cpp:5823 |
一个非贪心槽让整组退化为逐 token 走老路径 |
| ③ | 所有槽共用同一个 L、同一个窗口、同一个门铃 |
verify.cpp:1755-1827 |
严格 lockstep：慢槽拖住全组，没有"这轮这个槽先不算" |
| ④ | 每层主机屏障 + 16 个 staging 名额 | verify.cpp:1775-1811、verify.hpp:357 |
并发↑ ⇒ miss↑ ⇒ CPU 段↑ ⇒ 门铃等更久 |
② 值得单独说一句，因为它的影响面被严重低估：网页默认 temperature 0.6，不是贪心，
于是真实流量里几乎每个槽都带 sampled，整组直接不走链。
实测对照（lane A，同一模型同一负载）：
| 路径 | 1 槽 | 2 槽 | 3 槽 | 4 槽 |
|---|---|---|---|---|
| chain（贪心，走投机） | 194 | 121 | 82 | 67 tok/s·每请求 |
| batch（采样，逐 token） | 82 | 71 | 18 | 25 tok/s·每请求 |
同一份代码，只因为采样参数不同，单请求差 2.4 倍。 这条在 19.1 #3 里只是"惩罚类参数不生效"，
其实它连投机一起关掉了。
20.8 同一批数字的两种口径
改造时最容易吵架的地方：同一个系统，两套数字，差 2 倍。
| 口径 | 来源 | 典型值（4 槽） | 什么时候用 |
|---|---|---|---|
| 逐窗口统计 | STRATA_STEP_STATS=<path> → ~/logs/step-stats-a.jsonl（一窗一行 JSON） |
chain 4 槽：窗口 30.0 ms、行 8、接受 3.13 ⇒ ~138 tok/s·每请求 | 判"窗口机制快不快" |
| 压测台客户端 | ~/bench，流式时间戳 |
chain 4 槽：67 tok/s·每请求 | 判"用户实际感受到多快" |
差在哪：step-stats 的窗口是"混合了短上下文轻窗口"的中位数，而压测台是持续负载下的稳态。
两者都不假，但回答的不是同一个问题。
纪律（沿用 13 章）：性能数字必须带口径。 报"4 路 138 tok/s"而不说是逐窗口口径，
等于报了一个用户永远看不到的数。
step-stats 的字段（generate.cpp 的 step_stats.window(...)）：
{t, kind: batch|chain|prefill, slots, rows, drafts, accepted, step_ms, keeps, ctx}。
drafts/accepted 在这条 lane 上是判定"投机到底生效没生效"的唯一现成证据。
20.9 本章小结
- 一帧解码 92.5 ms 里，76% 是主机在算 CPU 专家、23% 是 GPU 在门铃上空转。GPU 不缺算力，缺的是活。
- 机制是每层一次门铃握手（
verify.cpp:1755-1827+verify_kernels.cu:504的单线程等待核），48 层 = 48 次。 strata batch:那行的per layer-window三项是计数不是毫秒；batch路径上三者之和 ≈ 行数 × 10，
chain路径的avg rows是 token 数、不能这么对账。- 三个硬上限里，
kStagingBlobs = 16是最便宜的：并发一高，超出的未命中全部掉进最慢的 CPU 档；32 只要 +30 MB。 - 并发退化的四个原因可以指认到四行代码；其中 ②「一个采样槽让整组退化」 用改一行默认值就能消除（194 vs 82）。
- 数字必须带口径：逐窗口统计与压测台客户端可以差 2 倍，回答的是不同问题。
下一章：21 改造手册：M0–M4 —— 知道病在哪之后，按什么顺序动刀。
本章出处
| 内容 | 出处（2026-10-21 在 SSEStraBed 上核过） |
|---|---|
| 92.50 / 70.00 / 21.14 分解 | ~/Strata/strata-iq3_xxs-a.log:200（2026-10-17 的历史日志，grep -n "strata batch:" 可复现；同文件 :223 是另一条 32.69 ms 的样本） |
| 汇总行格式与打印时机 | generate.cpp:5711-5725；batch_on() 在 :5615 |
| 门铃循环 | verify.cpp:1755-1827（pool(...) 在 :1797，*flag = want 在 :1811） |
| 单线程等待核 | src/kernels/cuda/verify_kernels.cu:504，启动点 :616 |
kStagingBlobs = 16 |
include/strata/core/verify.hpp:357；分配 verify.cpp:336，上限 :365 |
三级决策与 dma_src[64] |
src/core/expert_source.cpp:3376-3413 |
PCIe 实测与 pcie_frac |
lane A 日志 PCIe probe: 22.7 GB/s → pcie_frac 0.55；generate.cpp:413/2064-2077 |
L_max / sampled |
generate.cpp:5781 / :5823 |
| chain vs batch 每请求吞吐 | ~/bench 压测台，lane A，2026-10-21 |
| step-stats 字段与数值 | ~/logs/step-stats-a.jsonl（8014 batch + 4779 chain + 1102 prefill 条） |
| 四个硬上限汇总 | 10.4 与本章 20.6；--batch 解析 generate.cpp:1428，范围检查 :2960-2961（2…kVerifyMaxT） |
| step-stats 两个出口 | generate.cpp:5706（batch/chain 窗口）、:6600（prefill） |