第 9 章 prefill：最大的瓶颈

第 9 章 prefill：最大的瓶颈
MoE 大模型低成本部署框架第 9 章 prefill：最大的瓶颈
CHAPTER 09

第 9 章 prefill：最大的瓶颈

11 分钟阅读 · 2,844 字

问题清单 #1。本章回答三件事：为什么 prefill 是这套系统的头号瓶颈、
串行点具体在哪几行、已经存在与还缺的解法各自的代价。
数字出处标注日期；凡「需停服」的方案，落地状态以 docs/项目背景与问题总览.md 为准。

9.1 先把账算出来：为什么是它

硬约束 7（02 章）已经把结论写死：GPU 是唯一硬瓶颈，解码速度几乎与上下文长度无关，代价全在 prefill。
可观测性实验报告（2026-10-14 前后）把这组曲线钉死了：

量 实测 随上下文增长
解码速度 1k→258k：106→88 tok/s（−17%） 几乎不动
prefill 吞吐 ~3,350 tok/s 线性：TTFT ≈ 0.4s + tokens/3350
258K 提示词 ≈ 82s 一个请求独占 GPU 一分多钟

多用户并行的现实是：agent 形态的会话每一轮都携带全部历史重发（3.1 章书稿会话 = 每轮 15 万 token 起步）。
每个请求的 82 秒 × 请求数就是排队的期望长度。实测 8 路并发长上下文时 TTFT 可达 4.7 分钟
（问题清单 #1，2026-10-13 口径）；4 路 64k 并发时更狠——预填充进行时解码被饿死到 0.3~1.4 tok/s
（实验报告 §7.2）：在读新提示词的 GPU 周期里，所有人都不配说话。

一句话：decode 是"每人一小口"，prefill 是"一人一大口"。排队问题本质是 prefill 问题。

9.2 串行点在哪：两层三处

第一层：serve 的 FIFO（Python，最容易动）
Strata/serve/server.py 的 Service.fifo = threading.Lock()（:1268）：一个请求从进
prepare() 到生成结束整段持锁，别的请求连 tokenize 都不做（排队不暴露，观测盲区 §6.2）。
"一次只准一个请求读提示词"的直接实现就是它。
副作用：忙时新请求拿 ModelBusy 409（:125），而不是排队等待——客户端要自己重试。

第二层：引擎内部（C++，要重编 + 停服）

  1. src/prefill/prefill.cpp 约 :232-320：任务提交与 worker 由 std::mutex mu（:242）串行化，
    环（ring）staging 缓冲按序 DMA（:292/:334/:346 的自旋等待）。
    一次只有一个 prompt 的读取流水线在跑；文件头注释列出的优化路径
    （pinned arena、环形 staging、PeerPrefill、KV streaming）都是在"单流水线"内部做吞吐，不是并行化多 prompt。
  2. generate.cpp 的批步循环：prefill 与 decode 共享同一条 GPU 时间线，靠让出而不是交错。

9.3 已有的缓解：分段让出与检查点

① prefill 分段让出（#prefill-yield 段，generate.cpp 约 :6681-6692）：

STRATA_PREFILL_YIELD_CHUNKS = 每段读多少个批量 chunk（0 = 让出前行为：一口气读完）
STRATA_PREFILL_STEPS        = 每段之间让出多少个批步窗口给解码

语义是"读一段、让别人说几句、再读下一段"——饿死解码的问题由此缓和，
但总 prefill 时长不变（GPU 时间还是那些），TTFT 只是从"一个人 82 秒"摊成"人人慢一点"。

② 前缀缓存与会话停车（第 12 章展开）：重复历史不重读。这是唯一能把 82 秒变成 0.79 秒的机制
（64k 前缀命中实测，实验报告）。旋钮在 generate.cpp 的参数表（约 :439-447）：

--prompt-cache N          请求之间保留几个对话检查点（每个 ~118 MB 主机 RAM）——注释自称"TTFT 的关键旋钮"
--prompt-cache-every N    每 N 个新提示词 token 存一次检查点（默认 16384，0=关）
--prompt-cache-root N     根检查点密度（默认 2048）

③ 检查点为什么经常救不了（实测事故，2026-10-13）：会话停车快照单个 ~4.7G，而
--conversation-cache-mib 8192 的缓存装不下几个——全天驱逐 2508 次；被驱逐的会话下一轮
整段重读（八并发 TTFT 两分钟的主因之一，见 docs/reports 与 memory strata-ttft-prefill-queue）。
"每轮重读 238k"意味着缓存设计对大快照失效：密度（every=16384）太稀，重读区间大；
缓存容量又太小，存不下。→ 路线图旋钮：conversation-cache-mib 8192→24000+、
prompt-cache-every 16384→4096（路线图文档 §7.5，需停服窗口，截至 2026-10-19 未实施）。

9.4 四道 TTFT 治理旋钮（路线图第 1 项，均未实施）

旋钮 改动 代价 状态
停车缓存扩容 --conversation-cache-mib 8192 → 24000+ 主机 RAM（已 170+/216 GiB，内存悬顶） 待停服窗口
检查点加密 --prompt-cache-every 16384 → 4096 每检查点 ~118MB × 更多个 同上
普道准入 按提示词长度拒绝/延迟进入普道 策略复杂化、客户端体验 同上
组装预算封顶 RAG 时代的设计（预算 100-120k） RAG 已关，当前不适用 随 RAG 复活再议

9.5 根治方向：可并行 / 可恢复的 prefill（改造主线）

路线图第 2 项的引擎侧目标，拆开是两个词：

可并行：多请求的 prefill 同时段进 GPU（批 prefill）。当前内核的门槛：
专家分组 kernel 的每组条目上限 GMAX=8（s2_expert_grouped.cu:1360，
static_assert(GMAX >= kVerifyMaxT)；iq_kernels.cu:1804 的 kVerifyMaxT = 8 与 MMVQ
ncols<=8 模板硬编码）——批窗口超过 8 路，条目会被静默丢弃。
prefill 的"批"比 decode 的"批"大得多（每步几千 token × 多请求），撞的是同一组天花板。
这是 8 路以上天花板的物理来源（10 章再论）。

可恢复：prefill 中途可被打断、换别的请求用完 GPU、再回来续读。
分段让出（9.3①）已经有了"打断"，缺的是"续读不重读"——即检查点必须做到每段都可恢复
（prompt-cache-every 调密是第一步，恢复协议本身在改造计划里）。
docs/plans/ENGINE-BATCH-MTP-PLAN.md 与路线图把它列为"源码级重构"。

明确不做的：抢占式高速通道——代价是在途请求被取消、客户端要重试（已否决清单，02/19 章）；
最终选的是插队式（忙时新请求降级为批速度，不取消任何人）。

9.6 观测缺口：现在连"谁在排队"都看不见

问题清单 #1 为什么难治：总览 §6.2 的盲区表——prefill 排队详情（队列长度、谁在等、预计 TTFT）
不暴露。serve 的 FIFO 不记队列深度；/metrics 是单路状态；/v1/status 的 in_flight 是聚合数。
改造前必须先补观测（引擎侧输出队列/阶段 INFO 行即可），否则任何"并行化"都无法验证收益。

9.7 本章小结

  • 瓶颈的数学：TTFT ≈ 0.4s + tokens/3350 + FIFO 排队；排队项可以比前两项大一个数量级。
  • 串行点两层三处：serve FIFO（:1268）、引擎 prefill mutex（prefill.cpp :242 附近）、GPU 单时间线。
  • 缓解已有（分段让出、前缀检查点），根治未做（批 prefill、可恢复续读）；GMAX/kVerifyMaxT=8 是硬天花板。
  • 四个治理旋钮全部"已列未做"，卡在停服窗口与内存悬顶上。
  • 治它之前先看得见它：排队观测是改造的前置工作。

9.8 2026-10-21 复核增补（以本节为准，上文保留原样）

换到 SSEStraBed（8177 行）后逐行复核，§9.2 的两层三处有两处需要更正：

  1. serve 的主串行点不是 Service.fifo。生产 --batch 4，fifo 只在 --batch 0 时生效（server.py:1268/1825）。
    真正的串行是控制线锁 self.ctl：:657-660 注释原文「prompt reading, admission are taken one request at a time」，
    :670 取锁、覆盖 BGEN→BADM、:747 释放；BADM 之后 decode 才由 slot_q/BT 并发推进。
    因此「一次只准一个请求读提示词」在代码里仍然成立，只是实现从 fifo 换成了 ctl。
  2. 引擎侧没有 prefill 任务队列：generate.cpp:6075 是唯一的单线程 for (;;)，6831/6837 同步调用
    sp.run(...)；它返回前循环不会回到 6076 跑 decode。prefill.cpp:242 那个 mutex 是单流水线内部的串行，
    不是「多 prompt 排队」的实现。

新增事实（改造时必须知道）：

  • 粒度：--prefill auto 实际选 8192（generate.cpp:4426-4440，日志 prompt chunk auto: 8192 tokens）；
    #prefill-yield 的段 = STRATA_PREFILL_YIELD_CHUNKS(2) × chunk = 16384 token，段间跑 12 个窗口
    （6773-6791）。让出一次要 refill 专家缓存槽再重借，注释 ~300 ms（6770-6772）。
  • prefill 不走 CPU 专家池：prefill.cpp:1971-2034 的 fused 专家路径 + staging ring DMA，
    drive_pool 只在 decode/verify（generate.cpp:687-769）。所以 prefill 与 decode 争的是
    PCIe/DMA 与专家缓存槽，不是 CPU worker。
  • 借槽只在多 stage（--layer-split）发生：lend() 只处理 first >= 0 的 participant（6666），
    生产两条 lane 是单 stage，pf_parts[0].first = -1（4579）→ 不借槽、无 refill 成本，prefill 用自己的 scratch。
    （若将来开 layer split，才会出现「借 90%/85% 槽、host_res=kNotResident、refill ~300 ms」那套行为。）
    单卡上真正的耦合是专家缓存的驱逐竞争：prefill 每个 chunk 流式载入大量专家，会挤掉 decode 的热命中。
  • 不存在 batched prefill：sp.run 绑定单个会话 ss，in_lines 是无界 FIFO（:5418/5453/5460），
    BGEN 找不到空槽就直接报错（:6142-6155）。
  • GMAX/kVerifyMaxT=8 只约束 decode 的验证行数；prefill 走另一套 fused kernel，本轮没有找到
    「prefill 批撞同一组天花板」的证据（§9.5 该句待实验澄清）。
  • 路由：qwen3.8-flash→8092、-fast→8091 由 rag-proxy/config.json:68-84 按模型钉死，
    只有未命中才按 sha1(key) 粘性分流（rag_proxy.py:189-213）。

完整分析与参考方案：docs/reports/PREFILL-CONCURRENCY-STUDY-2026-10-21.md。

9.9 2026-10-21 复核增补之二：prefill 的专家路径有两条，差别巨大

§9.8 说"prefill 不走 CPU 专家池"，这句是对的；但它走哪条专家路径，直接决定 prefill 能不能被插队。
src/prefill/ 下是两条完全不同的实现：

路径 实现在 主机参与 生产是否在跑
MMQ（当前） prefill.cpp 的 mmq_plan()（:614-637） 每层一次全流同步 是
融合（fused int8） moe_fused.cu / moe_fused_iq.cu 专家分组在 GPU 侧完成，不逐层回主机 否（只差一个开关）

串行点就在这一行

// src/prefill/prefill.cpp:2055（MMQ 分支）
cudaStreamSynchronize(m.cs);    // 等待路由结果

它出现在 chunk 循环（:1629）→ 层循环（:1651）的最内层——每个 chunk 的每一层都要等一次全流。
prefill 的 chunk 是 8192 token、48 层，所以一次 sp.run 里会同步 48 次。
这不是"慢"，是这段时间内主机和 GPU 被绑死，任何解码窗口都插不进来（§9.8 的"不存在 batched prefill"由此而来）。

融合路径在本包上是可用的，只差 opt-in

三条都核过：

  1. 类型覆盖：moe_fused_iq.cu:1089 的 gu_covered = {IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S, IQ4_XS}，
    :1096 的 d_covered = {Q2_0, IQ4_NL}。
    本包（native_experts.txt 48 层逐行统计）：d 类型 30 层 Q2_0 + 18 层 IQ4_NL；
    gu 类型只有 IQ2_XXS(9) / IQ2_XS(10) / IQ2_S(10) / IQ3_XXS(6) / IQ3_S(13) 五种 ⇒ 全部在覆盖集内。
  2. 硬件门槛：moe_fused.cu:415 available() = dev_info().cc >= 80，4090D 是 8.9 ⇒ 满足。
  3. 开关：moe_fused.cu:428 的 requested() 要求显式 STRATA_PF_FUSED=1；生产 profile 没有设，
    日志里也没有 prompt experts on the fused int8 kernels 这一行。

⇒ 这不是"要改代码才能做的事"，是"没打开"。所以它排在改造手册 M0（21.2 ③）。

收益要诚实说：上游在原生包上的注释记的是 512 槽时相对 MMQ +8% / 0%（best/worst）。
真正的价值不是 prefill 变快，而是"prefill 不再每层卡主机"——那是 M3（decode 与 prefill 同窗）的前提。
另外它换了数值路径，必须核对同 prompt 贪心输出一致再上。

“用解码窗口读 prompt” 已经存在

§9.5 把"可恢复的 prefill"当改造目标。实际上代码里已经有原型：

// src/program/generate.cpp
int64_t short_read = 64;                                     // :459  默认只对 ≤64 token 的段启用
if (no_short || b - a > o.short_read) return false;          // :6547 门槛
auto read_windows = [&](int64_t a, int64_t b, ...) { ... };  // :6554 用解码窗口逐段读
...
sp_ok = read_windows(at, to, err);                           // :6826 调用点

也就是说："prompt 走 decode 窗口"这条机制已经写好了，只是被 64 token 的门槛限住，
用来处理"一条消息 / assistant 头 / 短工具结果"这类小段。
把它放开到 256/512 是零代码实验（21.2 ④），也是 M3 方向唯一能先验证收益的一步。

这条发现的意义在于改变了 M3 的起点：不是从零造一个"prefill/decode 同窗"，
而是把一条已经跑在生产上的旁路拓宽。 改造的难度评估因此下调一档。

全部章节