32 · 投机轮的锚定口径位级同流却输在一堵退不回去的墙上一句话把 MTP 接进投机轮次跑端到端 A/B——输出与纯自回归逐字节相同位级通过但加速比只有0.531×–0.737×全部 1×判否否因是结构性的被拒草稿行的 GDN 递推状态无法就地截断只能回滚后用M1 单步逐行重放而批量入口又硬要求批宽与mdl_mmax严格相等。「输出完全正确」与「能不能赚钱」之间隔着一堵结构性的墙。前置先读 第 31 篇MTP 上 GPU独立槽位、非致命绑定、逐级对拍。环境x86-64 · NVIDIA RTX 5090 · Qwen3.5-27Bq4 打包· CUDA 后端T64、k3、单会话。一、问题与结论第 31 篇把 MTP 头搬上了设备并且证明了它「抄对了」55 行逐级全过。但第 30 篇早就留了一句话「头能跑通 ≠ 头有用」。本篇就是去回答「有没有用」——结局是否定的而且否得很难看。指标数值判定判据 2与纯自回归同流T64 生成流memcmp 0六个起点全部成立✅ 位级通过实测判据 3端到端 ≥1.05×六个起点区间0.531×–0.737×❌ 判否实测每轮推进≈1.4 token/轮最好 1.730实测一轮成本≈35–40 ms实测理想草稿下E[len]2.107的推算上界≈0.76×推算一句话概括「输出完全正确」与「能不能赚钱」之间隔着一堵结构性的墙。二、背景先订正口径——什么叫「一轮」最初的写法是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) 草稿k3全从 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) verifyM k批的位置从 p1 起 mark() model_step_m(k, p1, [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_ii ≥ 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, p1, [d_1..d_{n_ok}]) 重放 n_ok 0 → rollback()再 model_step(1, p1, embed(自产 token))「不发出自产 token」那一条值得单独记它就是「位置pn_ok1的 token」而下一轮的d_1正是同一个位置的预测下一轮从pn_ok起锚。若这一轮先发下一轮接受d_1时会再发一次⇒ token 流出现重影。三、核心机制 / 做法3.1 要三个原语不是两个mark/rollback是一次性快照mark记下 GDN 的S/hist与各 attn 层的 KV 游标rollback把它们逐位恢复到mark那一刻然后快照失效必须重新mark。但投机轮次要的是一个中间值mark 在 p 快照点 消费到 p verify 跑到 pk 批已把游标推到 pk rollback → 回到 p 重放 n_ok 行 → 消费到 p n_ok ★ MTP 游标要停在 p n_ok —— 而快照里只有 p没有 pn_ok⇒ 必须新增vcuda_dev_mtp_truncate(dev, len)把 MTP 的写游标截到len不拷、不清 KV。/* 文件名GitHupSRC/src/npu/cuda/vllm_cuda_kernels.cu节选 *//* 把 MTP 那一层的**写游标截到 len**不拷 KV、不清 KV。 * 为什么需要它草稿链会把游标推到 pk而一轮只接受 n_ok 个 ⇒ 下一轮要从 * pn_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(len0||lens-mtp_maxseq)return-1;s-mtp_lenlen;s-mtp_trace_done0;/* 截断后快照不再配套 */if(s-mdl_mk_ok)s-mdl_mk_ok0;/* 在册快照一并作废 */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_snap0){s-mtp_lens-mtp_len_snap;s-mtp_trace_done0;}s-mdl_mk_ok0;/* 一次性回滚后快照失效必须重新 mark 才能再回滚 */GDN 那两笔cudaMemcpyS 149.6 MB hist就是本轮的贵处——见 3.5。3.2.1 回滚面要单独验一个负向对照专门验「MTP 游标是否真的进了快照面」臂做法结果正向reset → mark → 3 草稿 → rollback → 再跑同样 3 草稿out memcmp 0✅负向对照短路掉回退不动游标后重跑memcmp ≠ 0✅证明该面确实在起作用只跑正向是不够的如果回退本来就是空操作正向也会过。真正能证明「这一段在做事」的是把它拿掉以后结果变了。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_ok0 时是 SIZE_MAX * lgM SIZE_MAX*V 会回卷成 lgM - V ⇒ 在缓冲**前面**扫 993 KB访问违例。 * 而且语义上也对d_1 没被接受时 d_2 不可能被接受。 */while(n_ok1n_okkvamax_v(lgM(size_t)(n_ok-1)*V,V)d[n_ok1])n_ok;...if(n_okkv){/* 整批保留无需回滚 */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(intj1;jkv;j)seqB[emit]d[j];pkv;tkd[kv];}elseif(n_ok0){/* 只回滚不重放 */constintnxamax_v(lg,V);/* lg 未被动过 dist at p */if(vllm_cuda_model_rollback(cu)!0)gotodone;seqB[emit]nx;...}else{/* 重放 n_ok 行1..k-1。★ 只能用 M1 单步不能批量见 3.5。 */if(vllm_cuda_model_rollback(cu)!0)gotodone;for(intj1;jn_ok;j){if(vllm_cuda_model_step(cu,pj,hM(size_t)(j-1)*D,h_out,lg)!0)gotodone;seqB[emit]d[j];}pn_ok;tkd[n_ok];}3.4 拿到读数的实测日志[spec] 档A纯自回归,MTP064 token933.8 ms14.590 ms/token [spec] 档B投机,k3,force064 token1757.2 ms27.456 ms/token [spec] 轮数45 接受总数40 平均每轮接受0.889 ⇒ 每轮推进1.422 tokenn_ok0 的轮24 [spec] n_ok 直方图 n_ok0:24 n_ok1:6 n_ok2:11 n_ok3:4 [spec] token 流 memcmp0 ⇒ PASS与纯自回归**逐字节**相同★ 这条「构造保证」有一个代价第 33 篇会讲正因为发出流由主模型 argmax 兜底判据 2 对「喂给 MTP 的输入对不对」完全不敏感——它会在驱动的输入口径错得很离谱时照样报绿。3.5 那堵墙重放躲不掉因为批宽被钉死成本账一轮 ≈ 3 × c_d(≈3.4 ms) ← 三次设备 MTP 草稿步 3 行批(≈21 ms) ← verifyM3 n_ok × 14.6 ms ← ★ 重放被拒/接受前缀用 M1 逐步重放 ≈ 35–40 ms 一轮产出 ≈ 1.4 个 token ⇒ ≈ 23–27 ms/token vs 纯自回归 14.6那 21 ms 的 verify 批并不贵它是一次 M3 的批前向摊薄了权重读。贵的是重放n_ok k时每轮都要rollbackM1 逐行重放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||M2||Ms-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 判否起点首草稿命中率每轮推进档A档B加速046.7%1.422933.8 ms1757.2 ms0.531×2000052.4%1.524937.5 ms1571.1 ms0.597×5000043.9%1.561934.2 ms1495.3 ms0.625×10000056.8%1.730947.0 ms1284.6 ms0.737×最好15000043.2%1.455938.5 ms1665.1 ms0.564×24804563.2%1.684935.3 ms1520.1 ms0.615×全部 1×最好的数据也只有 0.737×离门槛 1.05× 差 1.4 倍。4.2 判据 2 位级通过T64 六个起点全memcmp0六起点与纯自回归的 token 流逐字节相同实测。它是构造保证见 3.3 实测坐实。4.3 归因订正一次「把推断当实测」的现形第一版的归因写的是「平均每轮接受 0.667≈ α₁(0.657)⇒第 1 个草稿之外几乎全落空n_ok0的轮占 62.5%」这句话有两个毛病①没实测0.667 ≈ 0.657是看着像不是量出来的②与工具自己的输出打架同一轮的工具自印轮数48 / 接受总数32 / n_ok0 的轮30⇒非零轮只有 18 个、平均接受 32/18 1.78——与「几乎全落空」直接矛盾。审计抓到了这条。补上n_ok直方图后重跑真相是n_ok0123轮数T6430783首草稿命中率 18/48 37.5%这才是低的那个量链尾并不落空条件命中d₂|d₁ 11/18 61.1%、d₃|d₁d₂ 3/11 27.3%平均0.667 32/48之所以低纯粹是被 30 个零接受轮拉下去的30 轮贡献 018 轮贡献 32。⇒与第一版归因正好相反。这条改写的价值不在于数字变了结论方向没变还是判否而在于它证明了这条纪律的必要性「先量后做」不是一句口号0.667 ≈ 0.657这种「数字碰巧接近」的推断一旦写进结论就会把「链尾失效」这个错误的机制解释固化下来误导后面的所有优化方向会去查「为什么d₂命中率低」而真问题在首草稿。五、边界与已知限制5.1 预登记的四条判据与处置开工前就冻结了判据与「判否后怎么办」#判据结果1位级·两端点强制全接受 / 全拒绝❌未单独跑——按原文字面产不出来见下2与纯自回归同流✅位级通过六起点memcmp03端到端 ≥1.05×❌判否0.531×–0.737×40逐字节❌ 未单独跑判据 1 为什么「字面产不出来」它要求「强制接受全部」这个端点。但实现里不能把n_ok硬置成k——若被接受的草稿本来就错硬置后整批的后续行是按错 token算的发出去的流就不再是主模型 argmax 链 ⇒ 判据 2与纯自回归同流会假红。所以--spec-force 1只做「观测 走状态路径」自然全接受时记成全接受并发往「整批保留、不回滚」那条状态路径没命中就按自然值走。⇒ 判据 1 记为待裁定项不事后改判据文字。5.2 判否后的动作预登记照办默认关保持VLLM_CUDA_Q35_MTP仍默认关主模型零接触不删代码--spec保留为已实现、可复跑的正确性载体它是目前唯一能在设备上位级证明锚定口径的装置记为「已实现但收益不足、不默认」。5.3 其余边界读数后来还被推翻过一次本篇的六起点读数是在「喂错 hidden」的真 bug 下测的修正后的幅度以 第 33 篇 为准判否方向不变。全部读数在同一会话、同一台机、真权重上取k3、T64。E[len] → 0.76×是推算不是实测它按离线E[len]2.107粗算重放项。判据 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 只退游标零拷贝批量入口硬要求Mmdl_mmax。真机验证命中E16MTP 接受率q35_mtp_accept ... --k 4α₁0.714p₁0.714、α₂0.364p₂0.254、α₃0.500p₃0.127、α₄0.750p₄0.095E[len]2.190 token/轮T63、teacher-forced 上界口径。锚定口径的投机--spec路径未独立复现读数后被第 33 篇推翻投机默认关仓内无 per-start 直方图落盘。六、小结「输出位级正确」与「有收益」是两条独立轴。这一轮把前者做到了逐字节后者仍然是 0.5–0.7×。任何「正确性通过 ⇒ 性能必然可期」的暗示都是错的。状态能否「退回去」是混合架构上投机的第一性约束。attn 的 KV 退游标是免费的GDN 的递推状态退回去要重算——而「重算」在批宽被钉死成一格时只能逐行来。⇒凡是「有递推状态」的层投机解码的收益都要先减掉「重放」这一项再谈。第三个原语的必要性来自「独立槽位」如果 MTP 的游标不是它自己的truncate这个操作根本没有定义第 31 篇的伏笔在这里收。构造保证会掩盖输入错误判据 2 之所以能位级通过恰恰因为它由主模型 argmax 兜底——这条性质让它对驱动的输入口径完全不敏感。第 33 篇就是踩在这上面。归因必须建立在分布上不是均值上。0.667这个均值同时兼容「每轮稳定 0.67」与「很多 0 少数 2」两种完全不同的机制只有直方图能分开它们。相关篇目第 31 篇MTP 上 GPU独立槽位 / 非致命绑定 / 逐级对拍第 30 篇MTP 模块与四个口径点第 19 篇 / 第 20 篇「游标即可见长度」的来源第 33 篇这一轮的读数为什么还不能作数。工单冻结、首试、验收与订正复测的细则见仓库GitHupSRC/docs/。源码与配套资源本仓库 https://gitee.com/pei-xiaoguang/kestrel-llm-cuda.gitCPU 推理源码 https://gitee.com/pei-xiaoguang/kestrel-llm
