Kestrel_VLLM - 32 · 投机轮的锚定口径：位级同流，却输在一堵退不回去的墙上

32 · 投机轮的锚定口径：位级同流，却输在一堵退不回去的墙上

一句话：把 MTP 接进投机轮次跑端到端 A/B——输出与纯自回归逐字节相同（位级通过），但加速比只有0.531×–0.737×（全部 <1×，判否）；否因是结构性的：被拒草稿行的 GDN 递推状态无法就地截断，只能回滚后用M=1 单步逐行重放，而批量入口又硬要求批宽与mdl_mmax严格相等。「输出完全正确」与「能不能赚钱」之间，隔着一堵结构性的墙。
前置：先读 第 31 篇（MTP 上 GPU：独立槽位、非致命绑定、逐级对拍）。
环境：x86-64 · NVIDIA RTX 5090 · Qwen3.5-27B（q4 打包）· CUDA 后端；T=64、k=3、单会话。

一、问题与结论

第 31 篇把 MTP 头搬上了设备，并且证明了它「抄对了」（55 行逐级全过）。但第 30 篇早就留了一句话：「头能跑通 ≠ 头有用」。本篇就是去回答「有没有用」——结局是否定的，而且否得很难看。

一句话概括：「输出完全正确」与「能不能赚钱」之间，隔着一堵结构性的墙。

二、背景：先订正口径——什么叫「一轮」

最初的写法是verify : model_step_m(4, pos, [t0,t1,t2,t3])——把「位置pos的 tokent0」和「位置pos的 hidden」配成一对。

但按第 30 篇那条口径：MTP 的输入是(embed(位置 i 的 token), 位置 i 的 hidden)，两者同位置配对。而一轮投机跑完之后，主模型手里只有**「最后被接受的那个位置」的 hidden——t0那一行（位置pos）早在上一步就被消费过了，再放进 verify 批就是同一个位置消费两次**。

订正后的锚定口径：

不变量：主模型已消费到位置 p；手上有 h_p、t_p、logits_p（= 位置 p 之后那一 token 的分布） MTP 自己的 KV 游标也在 p 1) 草稿（k=3，全从 p 出发、链式）： e = embed(t_p), h = h_p, q = p for j in 1..k: mtp_step(q, e, h, logits_mtp) ← 位置参数 = q d_j = argmax(logits_mtp) if j < k: e = embed(d_j); h = mtp_get(OUT) ← 链式喂**自己的输出** q += 1 2) verify（M = k，批的位置从 p+1 起）： mark() model_step_m(k, p+1, [d_1..d_k]) → h_out[k], logits[k] 3) 接受（★ d_1 用**已有的** logits_p，不占 verify 行）： n_ok = 1 若 argmax(logits_p) == d_1，否则 0 再逐位比 argmax(logits[i-2]) == d_i（i ≥ 2）直到第一次不等 4) 发出（最容易做错的一处）： n_ok ≥ 1 → 发出 d_1..d_{n_ok}（**不**发出自产 token） n_ok == 0 → 发出 argmax(logits_p) 5) 状态收尾（消费量必须与本轮发出量**相等**）： n_ok == k → 原地保留 n_ok ≥ 1 → rollback()，再 model_step_m(n_ok, p+1, [d_1..d_{n_ok}]) 重放 n_ok == 0 → rollback()，再 model_step(1, p+1, embed(自产 token))

「不发出自产 token」那一条值得单独记：它就是「位置p+n_ok+1的 token」，而下一轮的d'_1正是同一个位置的预测（下一轮从p+n_ok起锚）。若这一轮先发，下一轮接受d'_1时会再发一次⇒ token 流出现重影。

三、核心机制 / 做法

3.1 要三个原语，不是两个

mark/rollback是一次性快照：mark记下 GDN 的S/hist与各 attn 层的 KV 游标，rollback把它们逐位恢复到mark那一刻，然后快照失效（必须重新mark）。

但投机轮次要的是一个中间值：

mark 在 p （快照点 = 消费到 p） verify 跑到 p+k （批已把游标推到 p+k） rollback → 回到 p 重放 n_ok 行 → 消费到 p + n_ok ★ MTP 游标要停在 p + n_ok —— 而快照里只有 p，没有 p+n_ok

⇒ 必须新增vcuda_dev_mtp_truncate(dev, len)：把 MTP 的写游标截到len，不拷、不清 KV。

/* 文件名：GitHupSRC/src/npu/cuda/vllm_cuda_kernels.cu（节选） *//* 把 MTP 那一层的**写游标截到 `len`**（不拷 KV、不清 KV）。 * 为什么需要它：草稿链会把游标推到 `p+k`，而一轮只接受 `n_ok` 个 ⇒ 下一轮要从 * `p+n_ok` 起锚。`mark`/`rollback` 只能把它退回**快照那一格**，退不到中间值。 * ★ 为什么「只截游标」是安全的：保留下来的 [0, len) 全是**已确认**的 token * （可见长度/追加行号都取自游标，越界行永远读不到，下一轮 append 就地覆盖）。 */VCUDA_APIintvcuda_dev_mtp_truncate(void*dev,intlen){vc_dev_t*s=(vc_dev_t*)dev;if(!s||!s->ok||!s->mtp_ok)return-1;if(len<0||len>s->mtp_maxseq)return-1;s->mtp_len=len;s->mtp_trace_done=0;/* 截断后快照不再配套 */if(s->mdl_mk_ok)s->mdl_mk_ok=0;/* 在册快照一并作废 */return0;}

为什么可以不拷：MTP 那层 KV 的可见长度是由它自己的游标派生的⇒ 游标一改小，多写的那几行就永远读不到（下次 append 会就地覆盖）。所以「截断」= 退游标，零拷贝。这条论证与第 19/20 篇（KV 镜像的「游标即可见长度」）是同一条，只是这次它作用在一个自己的槽位上——这只有在第 31 篇「MTP 是独立槽位」成立时才成立。

3.2mark/rollback：两种状态，两种退法

同样的零拷贝论证也用在 attn 的 KV 上；但 GDN 的递推状态S/hist是就地覆写的，只能快照/恢复：

/* 文件名：GitHupSRC/src/npu/cuda/vllm_cuda_kernels.cu（节选：vcuda_dev_model_rollback） */if(cudaMemcpy(s->d_gdn_S,s->d_mb_ss,ssz*4,cudaMemcpyDeviceToDevice)!=cudaSuccess||cudaMemcpy(s->d_gdn_hist,s->d_mb_hs,hsz*4,cudaMemcpyDeviceToDevice)!=cudaSuccess){cudaGetLastError();return-1;}memcpy(s->attn_len,s->mdl_len_snap,(size_t)s->attn_nl*sizeof(int));/* 游标回退，零拷贝 *//* MTP 游标同样回退（`-1` = mark 时就没有 MTP 槽位）。 */if(s->mtp_len_snap>=0){s->mtp_len=s->mtp_len_snap;s->mtp_trace_done=0;}s->mdl_mk_ok=0;/* 一次性：回滚后快照失效，必须重新 mark 才能再回滚 */

GDN 那两笔cudaMemcpy（S= 149.6 MB +hist）就是本轮的贵处——见 3.5。

3.2.1 回滚面要单独验：一个负向对照

专门验「MTP 游标是否真的进了快照面」：

只跑正向是不够的：如果回退本来就是空操作，正向也会过。真正能证明「这一段在做事」的，是把它拿掉以后结果变了。

3.3 判据 2 位级通过：为什么它是「构造保证」

判据 2 要的是「端到端生成流 == 纯自回归参考流」。它成立的原因是构造性的：

• n_ok ≥ 1时，d_1..d_{n_ok}按定义就是逐位等于argmax(logits)（否则不会被接受）；
• n_ok == 0时，发出的是argmax(logits_p)本身。

⇒ 每一轮发出的 token都是主模型自己那条 argmax 链上的值，与草稿质量无关。代码里这条「接受 + 重放」是这么落的：

/* 文件名：GitHupSRC/tools/q35_gpu_model.c（节选：spec_run 的接受与状态收尾） */intn_ok=(amax_v(lg,V)==d[1])?1:0;/* ★ `n_ok >= 1` 这一项不能少：`(size_t)(n_ok - 1)` 在 n_ok==0 时是 SIZE_MAX， * `lgM + SIZE_MAX*V` 会回卷成 `lgM - V` ⇒ 在缓冲**前面**扫 993 KB，访问违例。 * 而且语义上也对：d_1 没被接受时 d_2 不可能被接受。 */while(n_ok>=1&&n_ok<kv&&amax_v(lgM+(size_t)(n_ok-1)*V,V)==d[n_ok+1])n_ok++;...if(n_ok>=kv){/* 整批保留，无需回滚 */memcpy(h_out,hoM+(size_t)(kv-1)*D,(size_t)D*4);memcpy(lg,lgM+(size_t)(kv-1)*V,(size_t)V*4);for(intj=1;j<=kv;j++)seqB[emit++]=d[j];p+=kv;tk=d[kv];}elseif(n_ok==0){/* 只回滚，不重放 */constintnx=amax_v(lg,V);/* lg 未被动过 = dist at p */if(vllm_cuda_model_rollback(cu)!=0)gotodone;seqB[emit++]=nx;...}else{/* 重放 n_ok 行（1..k-1）。★ 只能用 M=1 单步，不能批量（见 3.5）。 */if(vllm_cuda_model_rollback(cu)!=0)gotodone;for(intj=1;j<=n_ok;j++){if(vllm_cuda_model_step(cu,p+j,hM+(size_t)(j-1)*D,h_out,lg)!=0)gotodone;seqB[emit++]=d[j];}p+=n_ok;tk=d[n_ok];}

3.4 拿到读数的实测日志

[spec] 档A（纯自回归,MTP=0）64 token：933.8 ms（14.590 ms/token） [spec] 档B（投机,k=3,force=0）64 token：1757.2 ms（27.456 ms/token） [spec] 轮数=45 接受总数=40 平均每轮接受=0.889 ⇒ 每轮推进=1.422 token（n_ok=0 的轮=24） [spec] n_ok 直方图： n_ok=0:24 n_ok=1:6 n_ok=2:11 n_ok=3:4 [spec] token 流 memcmp=0 ⇒ PASS（与纯自回归**逐字节**相同）

★ 这条「构造保证」有一个代价，第 33 篇会讲：正因为发出流由主模型 argmax 兜底，判据 2 对「喂给 MTP 的输入对不对」完全不敏感——它会在驱动的输入口径错得很离谱时，照样报绿。

3.5 那堵墙：重放躲不掉，因为批宽被钉死

成本账：

一轮 ≈ 3 × c_d(≈3.4 ms) ← 三次设备 MTP 草稿步 + 3 行批(≈21 ms) ← verify（M=3） + n_ok × 14.6 ms ← ★ 重放：被拒/接受前缀用 M=1 逐步重放 ≈ 35–40 ms 一轮产出 ≈ 1.4 个 token ⇒ ≈ 23–27 ms/token vs 纯自回归 14.6

那 21 ms 的 verify 批并不贵（它是一次 M=3 的批前向，摊薄了权重读）。贵的是重放：n_ok < k时每轮都要rollback+M=1 逐行重放max(n_ok,1)行。

为什么重放躲不掉：批量入口有一条硬预检——批宽M必须与mdl_mmax严格相等：

/* 文件名：GitHupSRC/src/npu/cuda/vllm_cuda_kernels.cu（节选：vc_mb_ensure） */staticintvc_mb_ensure(vc_dev_t*s,intM){if(!s||M<2||M>s->mdl_mmax)return-1;/* ★ 硬约束：M 必须在 [2, mdl_mmax] */if(s->d_mb)return0;.../* ★ 按 mdl_mmax（上限）分配，而不是本次的 M —— 行距是 mdl_mmax，必须容纳全部行。 */if(cudaMalloc((void**)&s->d_mb,per*(size_t)s->mdl_mmax*4)!=cudaSuccess){...}return0;}

mdl_mmax是固定一格（3）——而重放的行数n_ok是1或2，凑不出「非满格批」⇒ 直接拒。

顺带一个真实踩到的 bug：尾轮kv = min(k, T-emit) = 2也会撞上同一条 ⇒「非满格批」被拒。修法是尾段不投机，走普通单步（只影响最后 ≤k−1 个 token）。

3.6 结构性判断：换掉草稿也救不回来

把「首草稿命中率」修到离线水平（E[len] → 2.107）呢？——重放那一项线性随n_ok增长，仍然把加速压在≈0.76×（推算）。

⇒这不是 MTP 质量问题，是「GDN 递推状态退不回去」问题。attn 层的 KV 可以「退游标」（第 19/20 篇已收口）；但GDN 的递推状态S/hist是逐 token 用S ← f(S, x)叠出来的，退到中间某一步必须重新逐 token 叠——没有位级捷径。加上「批宽必须严格相等」，这条退路就彻底封死了。

一句话概括这一轮：在混合架构上，投机解码的拦路虎从「草稿准不准」换成了「状态退不退得回去」。

四、实测数据

4.1 六个起点的完整读数（判据 3 判否）

全部 <1×，最好的数据也只有 0.737×（离门槛 1.05× 差 1.4 倍）。

4.2 判据 2 位级通过：T=64 六个起点全memcmp=0

六起点与纯自回归的 token 流逐字节相同（实测）。它是构造保证（见 3.3）+ 实测坐实。

4.3 归因订正：一次「把推断当实测」的现形

第一版的归因写的是：

「平均每轮接受 = 0.667≈ α₁(0.657)⇒第 1 个草稿之外几乎全落空（n_ok=0的轮占 62.5%）」

这句话有两个毛病：①没实测：0.667 ≈ 0.657是看着像，不是量出来的；②与工具自己的输出打架：同一轮的工具自印轮数=48 / 接受总数=32 / n_ok=0 的轮=30⇒非零轮只有 18 个、平均接受 = 32/18 = 1.78——与「几乎全落空」直接矛盾。

审计抓到了这条。补上n_ok直方图后重跑，真相是：

• 首草稿命中率 = 18/48 = 37.5%（这才是低的那个量）；
• 链尾并不落空：条件命中d₂|d₁ = 11/18 = 61.1%、d₃|d₁d₂ = 3/11 = 27.3%；
• 平均0.667 = 32/48之所以低，纯粹是被 30 个零接受轮拉下去的（30 轮贡献 0，18 轮贡献 32）。

⇒与第一版归因正好相反。这条改写的价值不在于数字变了（结论方向没变，还是判否），而在于它证明了这条纪律的必要性：

「先量后做」不是一句口号：0.667 ≈ 0.657这种「数字碰巧接近」的推断，一旦写进结论，就会把「链尾失效」这个错误的机制解释固化下来，误导后面的所有优化方向（会去查「为什么d₂命中率低」，而真问题在首草稿）。

五、边界与已知限制

5.1 预登记的四条判据与处置

开工前就冻结了判据与「判否后怎么办」：

判据 1 为什么「字面产不出来」：它要求「强制接受全部」这个端点。但实现里不能把n_ok硬置成k——若被接受的草稿本来就错，硬置后整批的后续行是按错 token算的，发出去的流就不再是主模型 argmax 链 ⇒ 判据 2（与纯自回归同流）会假红。所以--spec-force 1只做「观测 + 走状态路径」：自然全接受时记成全接受并发往「整批保留、不回滚」那条状态路径，没命中就按自然值走。⇒ 判据 1 记为待裁定项，不事后改判据文字。

5.2 判否后的动作（预登记，照办）

• 默认关保持（VLLM_CUDA_Q35_MTP仍默认关，主模型零接触）；
• 不删代码：--spec保留为已实现、可复跑的正确性载体（它是目前唯一能在设备上位级证明锚定口径的装置），记为「已实现但收益不足、不默认」。

5.3 其余边界

1. 读数后来还被推翻过一次：本篇的六起点读数是在「喂错 hidden」的真 bug 下测的，修正后的幅度以 第 33 篇 为准（判否方向不变）。
2. 全部读数在同一会话、同一台机、真权重上取；k=3、T=64。
3. E[len] → 0.76×是推算，不是实测：它按离线E[len]=2.107粗算重放项。
4. 判据 1/4 未单独跑，是流程内的已知缺口。

CPU 对照（迁移前基线）

• CPU 参考：src/model/vllm_q35_mtp.c:q35_mtp_step、src/model/vllm_q35_model.c:q35_model_step / q35_model_prefill—— 锚定口径的参考语义。
• 迁移要点：设备侧新增第三原语truncate（只退 MTP 写游标、不拷不清 KV，依据「游标即可见长度」）；rollback对 GDNS/hist做 D2D 快照恢复、对 attn KV 只退游标零拷贝；批量入口硬要求M==mdl_mmax。
• 真机验证：命中E16（MTP 接受率q35_mtp_accept ... --k 4：α₁=0.714（p₁=0.714）、α₂=0.364（p₂=0.254）、α₃=0.500（p₃=0.127）、α₄=0.750（p₄=0.095），E[len]=2.190 token/轮；T=63、teacher-forced 上界口径）。锚定口径的投机--spec路径未独立复现：读数后被第 33 篇推翻；投机默认关；仓内无 per-start 直方图落盘。

六、小结

1. 「输出位级正确」与「有收益」是两条独立轴。这一轮把前者做到了逐字节，后者仍然是 0.5–0.7×。任何「正确性通过 ⇒ 性能必然可期」的暗示都是错的。
2. 状态能否「退回去」是混合架构上投机的第一性约束。attn 的 KV 退游标是免费的，GDN 的递推状态退回去要重算——而「重算」在批宽被钉死成一格时，只能逐行来。⇒凡是「有递推状态」的层，投机解码的收益都要先减掉「重放」这一项再谈。
3. 第三个原语的必要性来自「独立槽位」：如果 MTP 的游标不是它自己的，truncate这个操作根本没有定义（第 31 篇的伏笔在这里收）。
4. 构造保证会掩盖输入错误：判据 2 之所以能位级通过，恰恰因为它由主模型 argmax 兜底——这条性质让它对驱动的输入口径完全不敏感。第 33 篇就是踩在这上面。
5. 归因必须建立在分布上，不是均值上。0.667这个均值同时兼容「每轮稳定 0.67」与「很多 0 + 少数 2」两种完全不同的机制；只有直方图能分开它们。

相关篇目：第 31 篇（MTP 上 GPU，独立槽位 / 非致命绑定 / 逐级对拍）；第 30 篇（MTP 模块与四个口径点）；第 19 篇 / 第 20 篇（「游标即可见长度」的来源）；第 33 篇（这一轮的读数为什么还不能作数）。
工单冻结、首试、验收与订正复测的细则见仓库GitHupSRC/docs/。
源码与配套资源：本仓库 https://gitee.com/pei-xiaoguang/kestrel-llm-cuda.git；
CPU 推理源码 https://gitee.com/pei-xiaoguang/kestrel-llm