Kestrel_VLLM - 27 · 寄存器常驻与融合发射：两次「少发一次」

用户36005557900
0 阅读34分钟

27 · 寄存器常驻与融合发射：两次「少发一次」

一句话：在 26 篇量清的「发射空档」之外，本篇走两条更直接的线——一次从寄存器入手（把 GDN 递推状态 S 常驻寄存器，省掉每步的全局内存往返，端到端 1.148×），一次从融合入手（把 β/a 的两次 GEMV 合并成一次发射，端到端 1.0411×）；共同点不是削计算也不是加并行度，而是少发一次。 前置：建议先读 第 26 篇（发射空档的两个机制与两条前置门禁）。 环境：x86-64 + NVIDIA RTX 5090；模型 Qwen3.5-27B q4；构建 -O2 -fmad=false -arch=native(sm_120)。 三个数：寄存器刀端到端 18.83 → 16.41 ms/token = 1.148×（同 exe 同会话交错 A/B；第 3.5 节）；融合刀端到端 16.528 → 15.875 ms/token = 1.0411×（省 0.653 ms/token）（两个只差一处的 DLL 交错 6 轮；第 4.5 节）；nsys 逐核实测 k_gdn_delta_reg<128> 4.97 µs/次 / 0.239 ms/token（同核通用版归档 55.1 µs ⇒ 11.1×；第 3.6 节）。 怎么读：结论看［一、问题与结论］与［六、诚实的边界］（必读）；机制看［三、第一刀 · 寄存器常驻］与［四、第二刀 · 融合发射］（想深挖再看）；各节的「审计结论与订正」（如 3.8）与文末附录是过程与原始数据（可跳）。

第 26 篇用三步量清了「发射空档」，得到一条可引用的结论：每次发射的空档 = 主机侧 CPU ≈5 µs/次 与 设备地板 ≈1.1 µs/次，取大的那个；模型实测的平均 2.05 µs 是两者混合。 它同时留下两项天花板候选——CUDA Graph 与 融小核，并划了「① 与 ② 不可相加」等三条前置门禁。

本篇是那两项候选之外、更直接的一条线：少发一次。两刀，方向不同—— 一次从寄存器入手（把 GDN 递推状态 S 常驻寄存器，省掉 shared/global 往返）， 一次从融合入手（把 β/a 的两次 GEMV 合并成一次发射）。 共同点既不是「削计算」，也不是「加并行度」，而是减少发射 / 访存次数。

状态：✅ 已落地。

一、问题与结论

  1. 两处「少发一次」都落地了，方向互不相同：
    • 寄存器这一刀：把 GDN 递推状态整列 S（128 个 float）装进寄存器， 两个 128 步循环内不再每步访问一次全局内存的 S[dk*vd+dv] ⇒ 端到端 18.83 → 16.41 ms/token，1.148×。
    • 融合这一刀：把 β/a 的两次 GEMV 发射合并成一次 k_f32_gemv2 ⇒ 端到端 16.528 → 15.875 ms/token，1.0411×（省 0.653 ms/token）。
  2. 共同点不是「削计算」或「加并行度」，而是「少发一次」。两次改动都不改算术、 不改归约序、不改地址表达式（寄存器那刀明文如此）；收益来自减少访存/发射次数，不是从算力里挖。
  3. 两次都先过 M7 形状门禁：动手改内核之前先用纯读地板量「形状级收益上限」， 门禁 基线/地板 ≥ 1.20×。寄存器刀的形状级门禁本样本 4.63×（审计独立复现 5.16×/5.31×）； 融合刀的门禁分母（融合轴）2.63/2.46/2.54×。
  4. 两次都位级/数值零回归。寄存器刀四层全部通过（S 0/786432、o 0/6144、maxrel 0.000e+00）； 融合刀探针四个候选全 bit-identical，模型路径 h/hfinal/logits 与改前归档逐位逐项相同。
  5. 两次都用「同会话、交错 A/B」而不是跨构建对比。寄存器刀的 ref 三次极差 0.05%、 寄存器版 0.6%；融合刀组内离散 BASE 1.14%、FUSE 0.56% ⇒ 1.04× 远超噪声。 融合刀先得出过的 17.00 → 15.87 = 1.072× 是跨构建对比，已作废。
  6. 两处「发射次数受限」都是推断，不是实测。融合刀的归因被审计显式降级为推断—— 纯读地板只给上界，没有 graph / 空核对照，分不开「主机侧发射」与「设备固定成本」。

mermaid diagram

图 1：本篇两刀「少发一次」的结构——都不削计算、不加并行度，只减少发射/访存次数。上支（寄存器刀）：k_gdn_delta 的两个 128 步循环原先每步都要访问一次全局内存的 S[dk*vd+dv]（且是依赖链），改为把整列 S（128 个 float）装进寄存器 Sr[KDT]、末尾一次写回；算式与归约序逐字未改，端到端 18.83 → 16.41 ms/token（1.148×，同 exe 同会话交错 A/B）。下支（融合刀）：β/a 原为两次独立 GEMV 发射（rows=48, cols=5120），合成一次 k_f32_gemv2，端到端 16.528 → 15.875 ms/token（1.0411×，省 0.653 ms，两个只差一处的 DLL 交错 A/B）。两刀的位级安全都以 -O2 -fmad=false 且不启用浮点重结合为前提。

二、背景：为什么下一刀选这两处

第 26 篇收在一句话上，这句话直接约束了「怎么改」：

每次发射的空档 = 主机侧 CPU ≈5 µs/次 与 设备地板 ≈1.1 µs/次，取大的那个。 短核走主机那一支（~5 µs），长核走设备那一支（~1.1 µs）；模型实测的平均 2.05 µs 是两者混合。

它给出的两条推论与本篇两刀严格对应：

  • 融小核能省发次数，但只对短核那一支有效（把发射压到设备地板）——对应本篇的 融合刀；
  • 那笔「被 GPU 的 16.6 ms 藏住的、≈5.7 ms/token 的主机 CPU 发射时间」一旦内核变快就会浮上来 ——对应本篇的 寄存器刀（把核变快、把每步的全局内存往返消掉）。

换句话说，第 26 篇里「跑道内那 ≈5 µs」，少发一次就少付一次；不增并行度、不上 graph， 也能直接吃掉它的一部分。

2.1 寄存器刀的对象：k_gdn_delta 的两难

第 26 篇的 ncu 已经证明 k_gdn_delta 是并行度 / 延迟饥饿： 48 block × 128 线程、DRAM 3.04%、SM 1.19%、0.02 waves/SM。因此下了判决—— 降字节数没用，要走「把工作切给更多 block」。

但真到切的时候，第 26 篇的三条禁令把候选筛掉了一大半：

  • 沿 K（kd）拆：Σ t·sk / Σ t·sq 是归约 ⇒ 改归约序 ⇒ 触碰保序纪律。排除；
  • 沿 vd（输出列）切更多 block：数值安全（每列独立、顺序不变）， 但不缩短每线程 128 步的串行链；而 48 个 block 本来就落在一波里（48 ≪ 170 SM） ⇒ 墙钟时间不变。排除。

真正的病灶不在并行度，而在每一大步都要访问一次全局内存的 S[dk*vd+dv]： 两个循环各 128 步，每步 1 读 + 1 写，而且是依赖链（第二步用的是第一步写回的值）。 ⇒ 于是有了第三种走法：不加 block、不切 K，改为把 S 搬进寄存器。这就是寄存器刀。

2.2 融合刀的对象：β/a 的 k_f32_gemv

寄存器刀落地后，nsys 逐核把 GDN 侧的排序重排了：

GDN 侧最大的内核不再是 k_gdn_delta，而是 β/a 的 k_f32_gemv —— 288 次（96/步）、8.60 µs/次、0.825 ms/token（min 8.19 / max 9.95，极稳）； k_gdn_delta_reg 已降到 0.239 ms。

ncu 对它的读数同样极端：只有 6 个 block、DRAM 6.00%、SM 0.75%。 与 k_gdn_delta 同性质——不是带宽瓶颈，是短核 / 发射问题。

β/a 每次是两次独立发射（a、b 两张同形矩阵各一次），形状真值固定 （vllm_cuda_kernels.cu:4069-4088）：rows = vh = 48、cols = dim = 5120、 wt = 256, wpb = 8 ⇒ grid = (48+8-1)/8 = 6 / block = 256。 两张矩阵同形 ⇒ 可以一次发射算两张。这就是融合刀：把两次发射融合成一次。

三、第一刀 · 寄存器常驻：k_gdn_delta 的 S 装进寄存器

3.1 方向：为什么是「S 常驻寄存器」而不是「切块」

按第 2.1 节的筛选结论，方向确定为：

把整列 S（128 个 float）装进寄存器，两个循环内只剩共享内存的 sq[]/sk[]。

这条路的漂亮之处在于它绕开了禁令：不改并行结构，就不触碰归约序； 不做算术改写，就不会动到浮点结合。它只改一件事——S 存在哪。 「同一段算式，变量从 global/shared 换成 register」，这是它位级安全的结构性理由。

落实就是给 k_gdn_delta 加一个模板化的兄弟版：Sr[KDT] 常驻寄存器，两条循环的算式与运算顺序逐字未改：

/* 文件名：src/npu/cuda/gdn_delta_twin.inc —— S 常驻寄存器版（KDT 编译期常量）
 * 「只改 S 存在哪」：把两个 128 步循环里对全局内存 Sm[] 的读写换成寄存器数组 Sr[]，
 * 算式（α 缩放 → S·k → 秩一加 → S·q）与运算顺序逐字未改 ⇒ 位级安全。 */
template <int KDT>
__global__ void k_gdn_delta_reg(float *__restrict S, const float *__restrict q,
                                const float *__restrict k, const float *__restrict v,
                                const float *__restrict g, const float *__restrict beta,
                                float *__restrict o, int kh, int kd, int vh, int vd) {
    const int vi = blockIdx.x;
    const int dv = threadIdx.x;
    if (vi >= vh || kd != KDT || vd != KDT) return;   /* 只服务 kd==vd==KDT，其它形状显式回退 */
    extern __shared__ float sh[];
    float *sq = sh, *sk = sh + KDT;
    ...
    float *Sm = S + (size_t)vi * (size_t)KDT * (size_t)KDT;   /* [KDT][KDT] */
    float Sr[KDT];
    #pragma unroll
    for (int dk = 0; dk < KDT; dk++) Sr[dk] = Sm[(size_t)dk * KDT + dv];  /* 一次读进寄存器 */

    const float alpha = expf(g[vi]);
    const float bt    = beta[vi];
    const float vv    = v[(size_t)vi * KDT + dv];

    float kv = 0.0f;
    #pragma unroll
    for (int dk = 0; dk < KDT; dk++) {          /* ① α 缩放 + (S·k)：不再每步读全局 */
        const float t = Sr[dk] * alpha;
        Sr[dk] = t;
        kv += t * sk[dk];
    }
    const float tmp = (vv - kv) * bt;
    float oa = 0.0f;
    #pragma unroll
    for (int dk = 0; dk < KDT; dk++) {          /* ② 秩一加 + (S·q)：也不再每步读全局 */
        const float t = Sr[dk] + sk[dk] * tmp;
        Sr[dk] = t;
        oa += t * sq[dk];
    }
    #pragma unroll
    for (int dk = 0; dk < KDT; dk++) Sm[(size_t)dk * KDT + dv] = Sr[dk];  /* 末尾一次写回 */
    o[(size_t)vi * KDT + dv] = oa;
}

3.2 新工具 tools/q35_gdn_shape.cu：M7 形状探针的方法论

M7 门禁的方法论只有一句话：先量形状级收益，够大才落地。 改内核之前先写一个纯读地板，把「这个形状最多能拿回多少」量出来；门禁判据 基线/地板 ≥ 1.20×。 不够大就不动手——这是把「改动成本」和「收益上限」对齐的护栏。

探针四档，形状逐项取自引擎真值（vh=48, kd=128, vd=128, kh=16， grid=48、block=128、动态 smem=1024 B）：

档内容
k_gdn_delta与引擎同一份源：#include "../src/npu/cuda/gdn_delta_twin.inc"（审计的强制版）
k_gdn_delta_reg<KD>同上（同一 #include）；只改「S 存在哪」：float Sr[KD] 常驻寄存器
k_gdn_delta_read<KD>纯读地板（只读 S 一遍 + 写 o） ⇒ M7 判据的分母
k_gdn_delta_readwrite<KD>读 + 写回地板（无计算），把「写 S」的代价单独分出

两个工程细节值得记：

  • 构建用 build_cuda_tool.ps1 -Source tools/q35_gdn_shape.cu（-O2 -fmad=false，与 DLL 同口径）；
  • .cu 含中文注释 ⇒ 必须带 UTF-8 BOM（下一刀里还会被再度升级为「BOM + 字符串字面量一律 ASCII」）。

为什么坚持「与引擎同一份源」：审计强制把两个内核抽到 src/npu/cuda/gdn_delta_twin.inc，引擎与工具都 #include 同一文件。 这就消掉了「改一边忘一边、M7 探针对着旧内核量」这个失败模式—— 探针量的必须是你要落地的那个内核，否则门禁本身就成了摆设。

3.3 微基准（RTX 5090，iters=200，两次独立运行）

档µs/次（两次）×48 层 ms有效 GB/s流量 MB
基线（引擎逐字）31.816 / 31.5191.527 / 1.513395.5 / 399.212.58（2R+2W）
S 常驻寄存器7.780 / 7.3390.373 / 0.352808.7 / 857.26.29（1R+1W）
纯读地板（只读 S + 写 o）6.372 / 6.8080.306 / 0.327493.7 / 462.03.15
读+写回地板（无计算）6.017 / 6.4400.289 / 0.3091045.6 / 977.06.29
  • M7 门禁：潜在加速上限 = 基线/纯读地板 ≥ 1.20× ⇒ 通过。 本样本读 4.63×（第二次运行），审计独立复现 5.16× / 5.31×。 ⚠ 纯读档自述抖 ±7% ⇒ 4.63× 是单次样本、不是精确值，不可单独外引。
  • 寄存器版实测 = 4.29×（审计独立复现 3.81× / 4.28×），已贴到读写地板（7.34 µs vs 6.44 µs）。
  • ⚠ 一个反直觉现象：「纯读」比「读+写回」还慢——纯读档有一个 128 深的 s += 串行加链， 读+写回档每步独立、并行度更好；这不影响任何结论，但说明这两档本身都不纯粹。
  • ptxas 证据（-Xptxas -v）：k_gdn_delta_reg<128> 255 registers / 40 B spill （≈10 个 float 落在 local）；k_gdn_delta 仅 40 registers。 ⇒ 128 线程/块 × 255 reg ⇒ 2 块/SM（48 块只占 24 个 SM）。
  • ⚠ 两次运行里纯读档抖了 6.8%（6.372→6.808），基线只抖 0.9%：小档读数按 ±7% 看。

3.4 位级一致性（四层，全部通过）

层级判据结果
微基准 memcmp同一个 S 输入、通用版 vs 寄存器版各跑一遍S 0/786432、o 0/6144、maxrel 0.000e+00 ⇒ 逐位相同
q35_gdn_check（f32）对 oracle 的容差闸门PASS，worst max_rel 3.800e-07（与归档完全相同）；0/5120、0/25600 ×2 两条位级自检逐位相同
q35_gdn_check --q4量级闸门 + 位级自检PASS，maxrel 2.721e-02（闸门 1e-1）；0/25600 ×2 逐位相同
q35_gpu_model --layers 64端到端对参考h 2.2536e-06 / hfinal 2.0334e-06 / logits 3.3887e-06，argmax 5/5，序列 148398 8442 198 198 198 —— 与归档逐项相同

未受影响的回归（本轮重建后重跑，均与归档逐位一致）：q35_rq_check maxrel 2.372e-06 / RMSrel 1.203e-06 / 0/17408 逐位相同 / PASS；q35_attn_check worst max_rel 4.697e-07 / PASS。

★ 位级证书的前提（审计订正）：「逐位相同」以下述编译口径为条件 —— -O2 -fmad=false（禁止把 a*b+c 收缩成 FMA）且未启用浮点重结合（nvcc 默认不重结合； 重结合需 -use_fast_math）。两版每一步都是「独立舍入的乘 + 独立舍入的加」， 累加器 kv/oa 各只有一个、都按 dk = 0..127 严格左折叠 ⇒ 同序、同舍入点。 Sr[128] 那 40 B spill 落在每线程私有内存，仍是精确 float 存取，不引入额外舍入。

⇒ 若改用 -fmad=true 或开启快速数学，两版可能产生不同的 FMA 融合模式， 证书失效、必须重验。不要把「算式逐字相同 ⇒ 逐位相同」读成无条件成立。

3.5 端到端 A/B（同一个 exe、同一会话交错重复 3 次）

开关：VLLM_CUDA_GDN_DELTA=ref（精确匹配；见第 3.7 节的风险条目）强制回通用版， 出厂默认 = 寄存器版。命令：

q35_gpu_model <md> --vqf <md>\model_q4.vqf --ref <md>\_q4ref64 `
    --layers 64 --lib build-x64\vllm_cuda.dll --iters 20
档三次 ms/tokentok/s有效权重带宽
ref（通用版）18.839 / 18.829 / 18.83153.1764–765 GB/s
默认（S 常驻寄存器）16.355 / 16.417 / 16.45960.8–61.1875–880 GB/s
差−2.423 ms/token（−12.9%）+7.71.148×
  • ★ 这个差不是抖动：ref 三次极差 0.05%、寄存器版 0.6% ⇒ 2.42 ms 远在噪声之外。 （微基准 hot 列 13% 的抖动是热/冷列，不是本端到端计时器。）
  • ref 档复现了归档端到端 18.9 ms / 53 tok/s ⇒ 口径对得上。

逐层的同轮 A/B（q35_gdn_check --q4，设备内、单层、T=5）：

层ref默认比
层 1（q4 三投影）0.130 ms/token（= 归档值）0.1001.30×
层 0（f32 三投影）0.4100.3501.17×

⚠ 这两行是 clock() 量化产物（步长 10 µs），只能做同轮 A/B，不可当绝对值引用。 q4 层的改幅度（1.30×）小于 f32 层（1.17×）是因为 f32 层每 token 要读 442 MB 的 f32 权重、 把 ⑤（k_gdn_delta 那份）的占比压下去了。

3.6 归因（★ 残差已用 nsys 直接测掉）

  • 归因是直接的：A/B 的唯一变量就是这个内核的实现（同 exe、同权重、同路径） ⇒ 端到端省的 2.42 ms 归它。

  • ★ 审计后补做的 nsys 逐核（★ 环境坑：nsys 的 protobuf 拒绝非 UTF-8 的 working_directory，必须把 cwd 也切到 subst 的 ASCII 盘，只 subst 参数不够； 产物一律落 %TEMP%，跑完删）：

    内核实例数总时（3 步）每次×48 层 ms/token
    k_gdn_delta（通用版归档）144—55.1 µs2.646
    k_gdn_delta_reg<128>（本轮）144715,890 ns4.97 µs0.239

    ⇒ 该内核 55.1 → 4.97 µs（11.1×），省 2.41 ms/token；与端到端 A/B 的 2.42 ms 吻合到 0.4% ⇒ 残差消失，本节第一版的「未解释」标注已作废。 本节由此取代第 26 篇归因表里 k_gdn_delta 那一行（其余行不变）。

  • 为什么微基准只有 4.29×，而引擎里是 11.1×：微基准是 back-to-back（S 3 MB 常驻 L2、 显存系统空闲）⇒ 它的基线偏乐观。量化：微基准基线 31.5 µs vs 引擎基线 55.1 µs （1.75× 拥塞放大）；寄存器版只读一遍 S、对拥塞不敏感 ⇒ 微基准低估了收益。 ⇒ 微基准只用于定方向与出位级证据，不能当收益；收益以引擎 A/B + nsys 为准。

  • 顺带的新事实（供下一步排序）：GDN 侧最大的内核不再是 k_gdn_delta，而是 β/a 的 k_f32_gemv —— 288 次（96/步）、8.60 µs/次、0.825 ms/token（min 8.19 / max 9.95，极稳）； k_gdn_delta_reg 已降到 0.239 ms。同轮其它大项：k_rq_q4_gemv_f32act 1203 次 / 11.93 ms、k_rms_norm_unit 387 次 / 0.357 ms。

这条「顺带的新事实」就是下一页的目录：融合刀的对象由此确定。

3.7 已知边界 / 未做

  • 已补（审计后）：nsys 逐核复测已完成（第 3.6 节表）⇒ 该内核 4.97 µs/次 / 0.239 ms/token 是实测值，0.23 ms 那处反算标注已作废；第 26 篇归因表的 k_gdn_delta 行已由第 3.6 节取代。
  • 形状覆盖只有 kd == vd == 128（本模型唯一形状）；其它形状走通用版（显式回退， 不是静默退化）。kd != vd 或非 128 的组合没有被验证过。
  • 255 registers / 40 B spill 是 -O2 -fmad=false -arch=native(sm_120) 的结果； 换 -arch 或换优化级别必须重新看 ptxas（spill 一旦变大就没有这个收益）。
  • ★ 同源风险条目（审计要求从「记一笔」升级为显式风险 + 验收）： VLLM_CUDA_GDN_DELTA（已由 VLLM_CUDA_GDN_SHAPE 改名）在出厂 DLL 里 也是认的 —— 它不是诊断档专属。
    • 风险：操作者误置 VLLM_CUDA_GDN_DELTA=ref ⇒ 静默掉性能（第 3.5 节里 ≈1.15× 那部分）。
    • 不构成静默错：它只切到「慢但逐位相同」的另一种实作（第 3.4 节已证）， 输出与产物仍是正确的。
    • 验收：(a) 判据必须是精确 strcmp(e, "ref") == 0 —— 第一版用 e[0] == 'r'， 会把 reg / random / 任何 r* 一并误判成「关」，已订正； (b) 发布前必须跑一次默认档 A/B，确认默认确实走寄存器版（本条为流程验收项）。
  • 同语义只写一份（审计已强制落地）：两个内核抽到 src/npu/cuda/gdn_delta_twin.inc， 引擎与工具都 #include 同一文件 ⇒ 不再存在「改一边忘一边、M7 探针对着旧内核」的失败模式。
  • 仍未做：非空闲显存系统下的对照微基准（用来把「1.75× 拥塞放大」直接量出来，而不是靠两个 不同运行拼出来）；冷启动（首次 launch）、其它 max_seq/batch、kd 非 128 的模型。
  • 产物：build-x64\vllm_cuda.dll 1,718,272 → 1,753,088 B（+34.8 KB，模板实例化； 订正 1–3 重建后尺寸不变）。

3.8 审计结论与订正

审计独立重建 DLL、独立跑完四层位级与端到端 A/B（未修改 src/、未修改 docs/）。

#级别结论处置
A成立位级：审计独立跑出 S 0/786432、o 0/6144、maxrel 0.000e+00，与本节逐项相同保留
B成立五条回归（gdn f32 / gdn --q4 / gpu_model 64 / rq / attn）逐项复现本节归档值，非复述保留
C门禁成立 / 需订正（小）门禁通过；但审计实测 M7 = 5.16× / 5.31×，本节 4.63× 落在其 ±7% 抖动区里 ⇒ 4.63× 是单样本✅ 已改为「≥1.20×（本样本 4.63×，复现 5.16–5.31×）」
D成立A/B 独立交错 3 次：ref 18.829/18.821/18.844、默认 16.432/16.465/16.473 ⇒ 差 2.374 ms / 1.144×，与本节 2.423 ms / 1.148× 同向同量级保留（区间写进引用口径）
E允许留白残差（微基准 4.29× 只该省 2.03 ms vs 实测 2.42 ms）：因被归档的是实测值本身（非外推），且已标「未解释 + 机制是假设 + 须补 nsys」⇒ 可留白。审计另用 nsys 归档量做了量级校核：55.1 µs vs 微基准 31.6 µs（≈1.74× 拥塞放大）×48 ≈ 2.645 ms，寄存器版 ≈7.4 µs×48 ≈ 0.355 ms ⇒ 省 ≈2.29 ms，与实测 2.42 ms 接近✅ 保留留白；但不得因此把它改写成「已证实」
F成立0.23 ms 已明确标注为反算，满足「反算必须标明」保留
1需订正发射点判据 e[0]=='r' 与注释「只有 ref 才关」不符 ⇒ 任何 r* 值都会静默关掉寄存器版✅ 已改为 strcmp(e,"ref")==0（且顺手改名 VLLM_CUDA_GDN_DELTA）
2需订正该开关在出厂 DLL 即生效，「记一笔」不足以覆盖「误置即掉性能」✅ 已升级为第 3.7 节的显式风险条目 + 两条验收（精确匹配；发布前跑默认档 A/B）
3需订正「两份独立拷贝、须双向同步」无强制 ⇒ 改引擎忘改工具会让 M7/位检对着旧内核✅ 已强制：抽到 src/npu/cuda/gdn_delta_twin.inc，两边 #include（该文件带 BOM）
4需订正「算式逐字相同 ⇒ 逐位相同」缺编译口径前提✅ 已补（第 3.4 节末：-fmad=false + 不重结合；-fmad=true/快速数学下证书失效）
5需订正M7 = 4.63× 被当精确值✅ 已改为区间写法（见 C 行）
6成立本系列的保序纪律、口径纪律、第 26 篇的三条禁令均未被违反；没有任何一句把推断/外推/反算写成实测—
7建议OMP_NUM_THREADS / VLLM_THREADS 未显式设置（纯 GPU 路径影响极小）记入待办，本轮未改
8已确认无误无仓库残留（审计自行清理了 6 个 .lib/.exp 副产物与 2 个 TEMP 文件），GPU 已释放—

审计的引用裁定：允许对外引用 1.148×，但必须带口径与区间：

「RTX 5090、64 层 q4 decode、T=1、--iters 20、同一 exe 同会话交错 A/B实测 ≈1.15× （本轮 1.148×；审计独立复现 1.144×）；两实作位级相同；适用边界 = 仅 kd==vd==128、 -O2 -fmad=false -arch=native(sm_120)。」

禁止：① 与第 26 篇的 1.8 ms graph 乐观外插相加；② 把 0.23 ms（反算） 或 4.29×（微基准） 与 1.148× 并列为同口径收益；③ 叠加进「19.24 → 14.0 ms」那类上限估计。

★ 审计后补做（把 E/F 两行的留白关掉）：审计说「残差可留白、但不得写成已证实」—— 下一轮直接把 nsys 逐核补上了（第 3.6 节表）：k_gdn_delta_reg<128> 144 次 / 4.97 µs / 0.239 ms/token（实测），对应通用版归档 55.1 µs / 2.646 ms ⇒ 省 2.41 ms，与端到端 A/B 的 2.42 ms 吻合到 0.4%。因此：

  1. E 行的「留白」已由实测关闭（不是被论证掉，是被测掉）；
  2. F 行那个 0.23 ms 反算值作废，改为实测 0.239 ms；
  3. 机制假设（微基准 L2 常驻 ⇒ 基线偏乐观 1.75×）得到实测支持（55.1 / 31.5 = 1.75）。

审计订正后的复验：订正 1–3 落地后重建引擎（DLL 仍是 1,753,088 B） 并重跑：gdn_check f32 worst max_rel 3.800e-07、0/5120、0/25600 ×2 逐位不变； gpu_model --layers 64 精度四项逐项不变；改名后的开关 A/B = 18.872 → 16.500 ms/token （1.144×），与订正前一致 ⇒ 订正只改源组织与开关名，未改变数值与性能。

四、第二刀 · 融合发射：β/a 的 k_f32_gemv 两次并一次

4.1 对象与门禁要求

账上留下的条目：β/a 的 k_f32_gemv 在 nsys 里是 96 次/token（a、b 各 48 层）、 8.60 µs/次、合计 0.825 ms/token，是 k_gdn_delta（寄存器刀）落地后 GDN 侧最大的一核。 ncu：只有 6 个 block、DRAM 6.00%、SM 0.75%。

门禁规定：动手改内核之前必须先做 M7 纯读 shape 探针（基线/地板 ≥ 1.20×）。

形状真值（vllm_cuda_kernels.cu:4069-4088）：rows = vh = 48、cols = dim = 5120、 wt = 256, wpb = 8 ⇒ grid = (48+8-1)/8 = 6 / block = 256。

4.2 改动

#位置改动
1新增 src/npu/cuda/f32_gemv_twin.inc把 k_f32_gemv 原样抽出（算式逐字符未改），并新增融合内核 k_f32_gemv2（一次发射算两张同形矩阵）。带 UTF-8 BOM
2src/npu/cuda/vllm_cuda_kernels.cuk_f32_gemv 定义处改为 #include "f32_gemv_twin.inc"（同 gdn_delta_twin.inc 既有做法）；vc_gdn_step_dev 里 β/a 的两次发射改一次融合发射 k_f32_gemv2<<<2*vh, 32>>>
3新增 tools/q35_f32_gemv_shape.cuM7 探针：基线 = 引擎同源内核；4 个候选工作形状；4 档纯读地板；冷/热两口径；四条轴分开量（同形状 / ILP / 并行度 / 融合）

k_f32_gemv2 的每半行：warp 划分、列步长 32、5 级 __shfl_down_sync 归约树与 k_f32_gemv 逐字相同，只是 warp ≥ rows 的第二半映射到 W1/y1：

/* 文件名：src/npu/cuda/f32_gemv_twin.inc —— 融合版：一次发射算两张同形矩阵（引擎里就是 β/a）
 * 与 k_f32_gemv 的差别只在「warp ≥ rows 时取 W1/y1」；每半行的列步长(32)、
 * 5 级 __shfl_down_sync 归约树与单张版**逐字相同** ⇒ 逐行结果逐位不变。 */
__global__ void k_f32_gemv2(const float *__restrict W0, const float *__restrict W1,
                            const float *__restrict x,
                            float *__restrict y0, float *__restrict y1,
                            int rows, int cols) {
    const int warp = (int)((blockIdx.x * blockDim.x + threadIdx.x) >> 5);
    const int lane = (int)(threadIdx.x & 31u);
    const int second = (warp >= rows);
    const int r = second ? (warp - rows) : warp;
    if (r >= rows) return;
    const float *wr = (second ? W1 : W0) + (size_t)r * (size_t)cols;
    float acc = 0.0f;
    for (int c = lane; c < cols; c += 32) acc += wr[c] * x[c];
    for (int off = 16; off; off >>= 1)
        acc += __shfl_down_sync(0xffffffffu, acc, off);
    if (lane == 0) { if (second) y1[r] = acc; else y0[r] = acc; }
}

调用契约：grid*block/32 ≥ 2*rows（写少了会静默算不完，if (r>=rows) return 拦不住）。 这条契约会被一个真实的坑验证——它是本篇最值得讲的一个坑（见第 4.8 节）。

4.3 位级 / 数值（逐项与归档相同）

探针 memcmp（合成数据）：4 个候选全部 A neq 0/48 B neq 0/48 [bit-identical]。

模型路径（同一个 q35_gpu_model，两个只差这一处发射形态的 DLL 交替跑）：

档h(layers)hfinallogitsargmax
VLLM_GEMV_AB_BASELINE=1（未融合）2.2536e-062.0334e-063.3887e-065/5
默认（融合）2.2536e-062.0334e-063.3887e-065/5

⇒ 三项 maxrel 与改前归档逐位逐项相同（打印值到全部有效位），argmax 与序列 148398 8442 198 198 198 相同。位级证书的编译口径前提：-O2 -fmad=false 且未启用 浮点重结合（同第 3.4 节）；换 -fmad=true / --use_fast_math ⇒ 证书失效。

4.4 M7 门禁（★ 判据按轴）

命令：q35_f32_gemv_shape --iters 50 --nl 512。冷档工作集 960 MiB = 10.0× L2 （L2 = 96.0 MiB，取自 deviceProp.l2CacheSize；审计指出上一版 90 MiB ≤ L2、 量到的是 L2 速率，已订正）。3 次同会话重复：

逐档原始读数见文末 附录 · 表 A2。

轴基线/地板判读
经典 M7（同形状同发射数）1.20 / 1.21 / 1.26×恰好压线
纯 ILP（U=4）1.33 / 1.26 / 1.29×用来排除该轴（且位级不安全）
纯并行度（6→48 block，两发）1.33 / 1.26 / 1.38×用来排除该轴（候选实测仅 1.04×）
融合（本轮门禁分母）2.63 / 2.46 / 2.54× ≥ 1.20 ✅通过
候选实测（融合 96x32）2.10 / 2.02 / 2.03×拿到融合轴天花板（0.405）的约 80%

第 4 个样本（最终重建的探针 exe、--nl 512）：引擎形态 1.044、融合 96x32 0.513 ⇒ 2.03×；经典 M7 1.23×、门禁分母 2.52×、位级仍 A neq 0/48 B neq 0/48。

★ 门禁纪律（本轮审计订正）：基线/地板 ≥ 1.20× 的地板必须与「被改的机制」同轴。 纯 ILP / 纯并行度两行只用于排除那两条轴，不是本轮门禁的分母 —— 否则会被读成 「门禁形同虚设」。另：只有只保留 load 的档才配叫「纯读地板」；仍执行全部 FMA 的档 只能叫「ILP 变体」（上一版命名不实，已改）。

这一条判据按轴，是本刀门禁与寄存器刀门禁最大的方法论差别：寄存器刀的地板只有一种（纯读）， 而本刀有四条轴，必须挑与改动同轴的那条当地板——否则「门禁通过」就是自欺。 改的是「发射次数」，所以分母只能是融合轴（2.63/2.46/2.54×），不是 ILP、不是并行度。

4.5 端到端 A/B（★ 同会话、交错、两个 DLL 只差一处）

命令：q35_gpu_model <md> --vqf <md>\model_q4.vqf --ref <md>\_q4ref64 --layers 64 --lib <dll>， build-x64-base\vllm_cuda.dll（-CudaDefines VLLM_GEMV_AB_BASELINE=1）与 build-x64\vllm_cuda.dll 交替各 5 次：

逐轮原始读数见文末 附录 · 表 A1。

组内离散：BASE 1.14%、FUSE 0.56% ⇒ 1.04× 远超噪声。

★ 口径订正（必须保留）：本轮先得出过 17.00 → 15.87 = 1.072×，那是拿改前归档值 当基线（跨构建、仪器状态不同）。同会话交错 A/B 的真值是 1.0411×（该节注记处写作 1.0399×，指同一交错口径）；以交错值为准，1.072× 作废。 跨构建对比只能作方向性证据。

⚠ --test-q35-cuda 的「5 token 前向」（17.04 → 16.30–16.49 ms/token）n 太小， 只作辅证，不作收益证据。

这两行合起来，就是本篇最值得记的一条纪律：1.072× 和 1.0411× 都「对」—— 它们量的是两件不同的事。收益数只能出自同会话、交错、单变量（两个 DLL 只差一处）的 A/B； 拿归档值当基线，看起来更省事，实际把构建差异也算进了收益里。

4.6 零回归

检查结果
L4 三自检 --test-sparse / --test-l3 / --test-moe-eprc 0 / 0 / 1，与归档一致
q35_gdn_check（f32）PASS，worst max_rel 3.800e-07（与归档相同）
q35_rq_checkmaxrel 2.372e-06 / RMSrel 1.203e-06 / PASS
q35_attn_checkPASS，worst max_rel 4.697e-07
Qwen3-VL-2B --stream-test33.82 tok/s（改前 33.68）；词表仍 151669 / im_start=151644
模型数值两个 DLL 的 h/hfinal/logits/argmax 逐项相同（第 4.3 节）

产物（锚在被改的那个产物：内核 ⇒ DLL）：

文件大小时间SHA256
build-x64\vllm_cuda.dll1,761,280 B12:29:4753B284F4B3E97BE84D48BF789B4AFB9DD12909F016D7610FE5C8D961B3EA9161
build-x64-base\vllm_cuda.dll（对照档）1,761,280 B12:29:18C25C408387E0F6AEF2811C5D8AD330A62C9DBD395DA692835E084058EC1C40B8
build-x64\vllm_kestrel_x64.exe1,570,047 B12:30:119C734FC1B7291738E485F1CDF483BFD1E9B50104EB51C27D2CC4C014358B1AD9
tools\build\out\q35_f32_gemv_shape.exe（探针）726,528 B12:29:125E24C32423DF1CF4DB9CF7B0930BE9A2CF89516C46063267E4EF5667C9007B8A

DLL 由 1,753,088 → 1,761,280 B（+8,192 B，融合内核实例化）。 exe 本轮源码未动（尺寸仍 1,570,047 B）；其 SHA 因整体重建而变 ≠ 「内容变了」。

4.7 归因分级（★ 把「推断」降级）

结论等级依据
「把 β/a 的两发并成一次」有 2.02–2.10×（形状级）/ 1.0399×（端到端）实测第 4.4 / 4.5 节
该核「发射次数受限」，而不是计算/并行度受限推断（与第 26 篇「主机侧 CPU ≈5 µs/次」同向）纯读地板只给上界；没有 graph / 空核对照 ⇒ 分不开「主机侧发射」与「设备固定成本」

这张表是本刀的诚实底线：形状级 2.02–2.10× 与端到端 1.0399× 是实测， 但「为什么有效」——k_f32_gemv 到底是发射次数受限——只是推断。 证据链里缺的正是第 26 篇那台机器（stream − graph 判别实验）： 没有 graph / 空核对照，就分不开「主机侧发射成本」与「设备固定成本」。 本条不得升格为实测。

4.8 本轮踩到的两个坑

坑 1（本篇最值得讲的一个）：grid 除错了 blockDim。

第一版发射写成 <<<(2*vh+31)/32, 32>>> = 3 个块（blockDim 已经是 32，不该再除）， 只算了 3 行、其余行留陈旧值。--test-q35-cuda 当场抓到： h maxrel 3.9454e-01 FAIL、序列变成 148398 8442 4960 198 198。 修成 <<<2*vh, 32>>> 后逐项回到归档值。

/* 错：grid 再除一次 blockDim ⇒ 只有 3 个块，其余行留陈旧值 */
k_f32_gemv2<<<(2*vh + 31) / 32, 32>>>(...);

/* 对：blockDim 已是 32，grid = 2*vh 才够 96 个 warp */
k_f32_gemv2<<<2*vh, 32>>>(...);

⇒ if (r>=rows) return 拦不住「块数不够」。这是一个「宁可拒绝、不要静默错」的反例： 块数写少，边界检查反而把越界掩盖成「剩下的行保持原样」——静默错。 所以调用契约必须写死（grid*block/32 ≥ 2*rows），而不是靠内核里的守卫。

坑 2：tools/*.cu 的 cp936 吞引号。

新增 .cu 带中文字符串字面量时，cl.exe 按 cp936 解码、 中文尾字节与 " 配对把引号吞掉 ⇒ error: missing closing quote。 实测「补 UTF-8 BOM」并不充分（有 BOM 仍复现）⇒ 纪律升级为 「BOM + 字符串字面量一律 ASCII」（注释仍可中文）。

4.9 仍需注意 / 未做

  1. 编译期诊断档：-CudaDefines VLLM_GEMV_AB_BASELINE=1 可编出「未融合」的对照 DLL。 它是编译期开关（不是运行期 env）⇒ 出厂 DLL 里根本不含那条路径， 不会重演第 3.7 节那种「误置 env 静默掉性能」。但必须放另一个 -OutDir。
  2. 融合轴的绝对天花板只到 0.405 ms/48 层（纯读）。已落地档 0.507 ⇒ 还差 ~20%， 但那 20% 里含计算，位级安全的做法未知（只能沿 vd 切更多块，且不缩短串行链）。
  3. 未做：graph / 空核对照（把「主机侧发射」与「设备固定成本」分开）；本文的探针没有 测发射空档，故第 4.7 节的归因仍是推断。
  4. 未做：k_gdn_beta_g（48 线程、2 block）等其余 GDN/attn 小核的融合（第 26 篇记的 192 次发射，0.39–0.86 ms）—— 与本轮同轴（发射次数），是下一个自然候选。
  5. 环境：僵死 vllm_kestrel_x64 进程仍为 7 个，Stop-Process -Force 无效，需重启机器。

五、共同方法论：一条闭环

把两刀并排看，形状完全一样——它们不是两个技巧，是同一条流水线：

① M7 形状门禁（纯读地板，基线/地板 ≥ 1.20×）
     ↓ 够大才动手
② 落地位级 / 数值零回归（四层位检 / 逐项归档相同）
     ↓
③ 同会话、交错、单变量 A/B（同 exe 或两个只差一处的 DLL）
     ↓
④ 审计（独立重建、独立复跑）
     ↓
⑤ 归因分级（实测 / 推断 / 上界 / 作废，逐条标注）

逐步对照：

环节寄存器刀融合刀
① 形状门禁纯读地板 ≥1.20×（本样本 4.63×，复现 5.16–5.31×）按轴取融合轴分母 2.63/2.46/2.54×
② 位级四层全过（0/786432、0/6144、0.000e+00）四候选 bit-identical、三项 maxrel 逐项相同
③ A/B同 exe、同会话、交错 3 次 ⇒ 1.148×两个只差一处的 DLL、交错 6 轮 ⇒ 1.0411×
④ 审计独立重建、独立复跑 ⇒ 1.144× 复现零回归 + 归因降级
⑤ 归因残差被 nsys 直接测掉（2.41 vs 2.42 ms，0.4%）降级为推断（缺 graph/空核对照）

两刀给方法论各补了一条：

  • 寄存器刀教的是「探针必须与引擎同源」：gdn_delta_twin.inc 的强制抽取， 是把一个「纸面纪律」变成「编译期强制」的动作。探针量错内核，门禁就是摆设。
  • 融合刀教的是「门禁判据要按轴」：四条轴里只有「融合轴」是这一轮的分母， 另外两条（ILP、并行度）的作用是排除，不是分母。用错了轴，≥1.20× 就形同虚设。

而两刀共享的两条纪律，是从第 26 篇直接继承的：

  1. 收益只出自同会话交错 A/B——1.072×（跨构建）作废，1.0411×（交错）为真；
  2. 绝不能把推断写成实测——融合刀的归因被显式降级，寄存器刀的「发射次数」也只在 k_f32_gemv 那条链上有形状级上界，没有直接观测。

六、诚实的边界

本篇两刀的共同点，是一个很窄的命题：减少发射 / 访存次数能直接吃掉主机侧的 ≈5 µs。 但这句话对为什么有效的解释，两次都停在同一个位置：

「发射次数受限」是推断，不是实测。

  • 寄存器刀：残差被 nsys 测掉了（55.1 → 4.97 µs，与端到端 2.42 ms 吻合到 0.4%）， 这一条是实测；但「为什么微基准 4.29× 而引擎 11.1×」的机制（L2 常驻 ⇒ 基线偏乐观） 是实测支持下的假设（55.1 / 31.5 = 1.75），不是被单独量出来的。
  • 融合刀：形状级 2.02–2.10×、端到端 1.0399× 是实测；但「这一核发射次数受限， 不是计算/并行度受限」是推断——缺 graph / 空核对照， 与第 26 篇的判别实验同向，却没有那台机器去把「主机侧」和「设备侧」分开。

所以本篇的收益数字可以引用，但必须带来源与边界：

优化可引用收益来源 / 口径边界
寄存器刀（S 常驻寄存器）1.148×（复现 1.144×）同 exe 同会话交错 A/B仅 kd==vd==128、-O2 -fmad=false -arch=native(sm_120)
融合刀（融合两次发射）1.0411×（省 0.653 ms）同会话交错、两个只差一处的 DLL融合轴天花板 0.405；归因为推断

禁止（寄存器刀审计明列）：把 0.23 ms（反算） 或 4.29×（微基准） 与 1.148× 并列为 同口径收益；把寄存器刀与第 26 篇的 1.8 ms graph 乐观外插相加；叠加进「19.24 → 14.0 ms」 那类上限估计。

CPU 对照（迁移前基线）

  • CPU 参考：src/model/vllm_q35_gdn.c:q35_gdn_step —— S 递推与 β/a 的 CPU 参考。
  • 迁移要点：两处「少发一次」——① 把整列 S(128 float) 装进寄存器（k_gdn_delta_reg<KDT>，只改 S 存在哪、算式/归约序逐字未改）；② β/a 两次 GEMV 合成一次 k_f32_gemv2；均先过 M7 形状门禁、位级/数值零回归。
  • 真机验证：命中 E7（q35_gdn_check 覆盖：5/5 OK、host↔设备位级一致；k_gdn_delta_reg 为默认档）；融合刀端到端 1.0411×（同会话交错 A/B）与寄存器刀 1.148× 属同层内，未在提供证据内。

附录：原始数据表

表 A1 · 端到端 A/B：同会话交错 6 轮（原第 4.5 节）

轮BASE ms/tokenFUSE ms/token比值
116.54015.8431.0440×
216.44415.8541.0372×
316.52015.9311.0370×
4（重建后复验）16.57015.8831.0433×
5（重建后复验）16.46215.8541.0383×
6（最终产物复验）16.63015.8841.0470×
6 轮均值16.52815.8751.0411×（省 0.653 ms/token）

表 A2 · M7 门禁逐档读数（3 次同会话重复）（原第 4.4 节）

档（ms / 48 层）run1run2run3
引擎形态（A+B 两次发射）1.0631.0261.036
候选 融合 12x2560.5250.5050.510
候选 融合 48x640.536——
候选 融合 96x32（已落地）0.5070.5090.511
地板 纯读 U=1 @6x256（同形状同发射）0.8400.8510.863
地板 纯读 U=1 @48x32（并行度轴）0.7970.8140.750
地板 纯读 U=4 @6x256（ILP 轴）0.8000.8150.800
地板 纯读 U=1 融合 @96x32（融合轴）0.4050.4180.407

七、小结

  1. 「少发一次」是一条独立的优化轴：不削计算、不加并行度，只减少发射 / 访存次数， 就能吃掉主机侧那 ≈5 µs 的一部分。
  2. 寄存器刀绕开禁令靠的是「只改 S 存在哪」：不改并行结构、不改算术、不改归约序， 位级安全是结构性的——但位级证书仍以 -O2 -fmad=false 且不重结合为前提。
  3. 探针必须与引擎同源：把内核抽到 .inc、两边 #include，是「编译期强制」， 消掉「改一边忘一边、门禁对着旧内核」的失败模式。
  4. 门禁判据要按轴：改的是发射次数，分母就只能取「融合轴」；ILP / 并行度两轴只用于排除。
  5. 收益只能出自同会话、交错、单变量的 A/B：拿归档值当基线把构建差异也算进了收益 （1.072× vs 1.0411×）。
  6. 微基准只用于定方向与出位级证据：back-to-back 下基线偏乐观（L2 常驻，拥塞放大 1.75×）， 真收益以引擎 A/B + nsys 为准。
  7. 推断与实测要分级：融合刀的「发射次数受限」是推断（缺 graph / 空核对照）， 不得升格为实测。

相关篇目：第 26 篇、第 28 篇 源码与配套资源：本仓库 gitee.com/pei-xiaogua…； CPU 推理源码 gitee.com/pei-xiaogua…

avatar
文章
阅读
粉丝