Kestrel_VLLM - 34 · 批量 M 行：五个块、两种层型，与「退回去」的两个原语

用户36005557900
0 阅读18分钟

34 · 批量 M 行：五个块、两种层型，与「退回去」的两个原语

一句话：为了「一次前向验 k+1 个位置」，给 hybrid 设备路径造一条批量 M 行的路径——它必须与逐行单 token 逐位等价（不是「误差很小」）；五个块（a–e）、两种层型（GDN 与 full-attn 各有各的批量难点）、外加两个只为投机存在的原语（mark / rollback）。全部判据 memcmp = 0。 前置：建议先读 第 32 篇（它用的是这些基建）与 第 19 篇 / 第 20 篇（「游标即可见长度」的来源）。 环境：x86-64 · NVIDIA RTX 5090 · Qwen3.5-27B（q4 打包）· CUDA 后端；M ∈ {2..8}。

一、问题与结论

第 32 篇说「hybrid 设备路径是纯 M=1 的，没有批量 verify」——本篇讲把它建出来的全过程。这是阶段五工程量最大的一块。

指标数值 / 结果判定
要建的东西① 投影的 M 行；② GDN 的 M 维；③ full-attn 的 M 行；④ 状态克隆/回滚（配套）—
端到端等价（--mt M）h_out / logits M ∈ {2..8} 全 memcmp = 0实测（逐位）
内核级等价tools/q35_attn_m.cu 16/16 全 PASS（M ∈ {1,2,4,8} × base ∈ {0,1,5,17}）实测
回滚等价（--rb M）真实混合与合成档 M ∈ {2,4,8} 全 memcmp = 0实测（有区分度）
不启用就不变M=1 三文件 SHA256 / 尺寸不变实测

一句话概括：批量的每一步都被一条「必须逐位相等」的判据卡着——这条纪律在三次真实的翻车里救了场。

二、背景：先冻结，再落地

这一段和投机线的其它部分一样，先冻结后落地：

步骤内容为什么这样切
第一步把批量投影内核搬进引擎，只对 M=1 生效（M=1,R=1 应与现役逐位相同）它不改变任何行为 ⇒ 零风险、可独立提交
第二步打开 M>1：① 投影按形状选 R；② GDN 的 M 步串行；③ attention 的 M 行 + 块内因果；④ 状态克隆/回滚真正的工作量；默认关（env 门）
↳ 内部顺序先 --synthetic attn 落 ①③（完全绕开 GDN）⇒ 再 --synthetic gdn 落 ② ⇒ 最后 ④工具已经支持 --synthetic gdn|attn：把两种层型解耦成两次独立、各自可验证的落地
第三步接 spec_verify（草稿来自 MTP 头）最后一步；此时才谈端到端

第一步的一个漂亮性质：k_rq_q4_gemv_mr<R,M> 是模板，而本步没有任何调用点 ⇒ 一个实例都不会被实例化 ⇒ 不产生任何代码 ⇒ DLL/exe 尺寸不变。

这也提醒一条口径：PE 内嵌构建时间戳 ⇒ SHA 必变，所以这里用「尺寸相同 + 模板零实例化」作为证据，而不声称 SHA 相同。

2.1 R 表（调参面，不是常量）

R 是按形状选的，不是一个常数：

形状最优 R理由
mlp gate/up、gdn in_qkv、attn q、lm_head4R=4 明显更好
mlp down、gdn out / attn o2R=4 与 R=2 持平或略差
attn k/v、gdn in_proj_z2（甚至 1）R=4 是负的：行数少 ⇒ warp 数少 ⇒ 并行度不足

这张表落地时以实测为准重扫，不照抄。

三、核心机制 / 做法

3.1 块 a–e：五个块，一条判据

最小可验收单元不是「再加一个内核」，而是这一整块：

块内容位级判据
aM 行 RMSNorm / add / swigluM=1 与现役逐位相同
b投影批量（按 R 表）每行 == 单行版
cRoPE 表 M 行（行 m 用位置 pos+m）+ q/k prep M 行同上
dKV append M 行 + attention M 个 query 行（行 m 只看 prefix+m+1 个 key）同上
evc_model_step_body_m + env 门（默认关）开 M=1 时三文件 SHA256 不变

验收口径：M=4 的批量前向必须与同一段文本 4 次单 token 前向逐位相同——因为每行的算术序没有变，这应当是精确成立的，不是「误差很小」。

3.2 block a 的翻车：验证拿不到，就回退

block a 改完（k_rms_norm_rows 替换三处 norm 调用），DLL 尺寸从 1,783,808 → 1,800,192——因为它是非模板、未调用也发射代码 ⇒ 「尺寸相同」这条捷径失效，必须实测。可当轮拿不到实测，两个原因：

  1. 归档的 dump 哈希复现不了：归档是 T=64 真实文本，工具默认是合成 T=1 ⇒ 输入不同。规格里只留了哈希、没留「怎么跑出来」。
  2. A/B 的基线被自己覆盖了：改动前的 DLL 已被重建覆盖（旁路目录里是另一个档，不是「改动前」）。

⇒ 按「先验证后落地」回退，重建确认回到 1,783,808 B。

第二轮按流程护栏重来，成功：① 改前先把现役 DLL 复制到 build-x64-prev\；② 定一条规范的 dump 命令（补上「不可复现」的缺口）；③ 改动前的二进制 dump 三文件 SHA256；④⑤ 关/开改动各重建一次 dump ⇒ 三文件逐字节全等。

由此立了两条流程纪律：① 改引擎前先把现有 vllm_cuda.dll 复制到旁路目录；② dump 的哈希必须连「完整命令行 + 输入来源」一起归档 —— 只留哈希等于没留。

block a 另一条对后续通用的结论：k_add_ip / k_swiglu 不需要新内核——它们是纯按下标的逐元素内核，[M][d] 连续布局下直接传 n = M*d 就逐位等价。

3.3 block b：一行改动，换来「天然可寻址 M 行」

把投影发射器由 k_rq_q4_gemv_f32act 改成 k_rq_q4_gemv_mr<1,1>（一行）。意义是结构性的：投影发射器从此天然可寻址 M 行——M>1 只需把指针换成行主序 [M][·]、grid 换成 M，不必再动这条路径的结构。位级：三文件 SHA256 全等；热路径 +0.09% / +0.22%（噪声内）。

3.4 block c/d：M 的编码方式，与「块内因果」

M 的编码方式很讲究：不加内核参数，由 gridDim.x 派生：

内核grid行号 / 行内下标
k_attn_kv_appendM*bpr_kvm = blockIdx.x / bpr_kv；blk = blockIdx.x % bpr_kv；目的行 = dyn[1] + m
k_attn_decodeM*nhm = blockIdx.x / nh；i = blockIdx.x % nh

为什么不用内核参数：graph 捕获时参数被定死，而从 gridDim.x 派生与参数无关。这是第 29 篇「把随 token 变化的量赶出内核参数」的同一条思路。

★★ block d 的真正新东西是「块内因果」：k_attn_decode 的可见长度由 npos = dyn[1] + 1 变成 npos = dyn[1] + m + 1——第 m 个草稿行只能看到 prefix+m+1 个 key（它自己及之前的），看不到后面的草稿行。这正是「批量 verify 与逐行单 token 等价」的根据。M=1 时 m ≡ 0 ⇒ 与单 token 版同一表达式。

这段两内核被抽到唯一一份源 attn_core_twin.inc，引擎与工具 #include 同一文件 ⇒「改了引擎忘了改工具」这个失败模式被结构性消除：

/* 文件名：GitHupSRC/src/npu/cuda/attn_core_twin.inc（节选） */
/* ⑤ KV 追加：把 kt/vt 里的 **M 个** [kv_width] 行，写到 cache 的 `dyn[1] + m` 行。 */
__global__ void k_attn_kv_append(float *__restrict kc, float *__restrict vc,
                                 const float *__restrict kt,
                                 const float *__restrict vt,
                                 const int *__restrict dyn, int kv_width) {
    const int bpr = (kv_width + blockDim.x - 1) / blockDim.x;   /* 每行块数 */
    const int m   = blockIdx.x / bpr;
    const int blk = blockIdx.x - m * bpr;
    const int row = dyn[1] + m;
    const size_t off = (size_t)row * (size_t)kv_width;
    ...
}
/* GQA 注意力（一个 block 一个 (行 m, q-head i)）
 * ★ 可见长度 = dyn[1] + m + 1（**块内因果**：行 m 看不到后面的草稿行）。 */
__global__ void k_attn_decode(float *__restrict ao, const float *__restrict qg,
                              const float *__restrict kc, const float *__restrict vc,
                              const int *__restrict dyn, int nh, int nkv, int hd,
                              int kv_width, float scale) {
    const int m = blockIdx.x / nh;
    const int i = blockIdx.x - m * nh;
    const int d = threadIdx.x;
    if (i >= nh) return;
    const int npos = dyn[1] + m + 1;
    ...
    float *scp = sh;                    /* [npos] */
    float *red = sh + npos;             /* [blockDim] */

★ 动态 shared 的定尺是调用方契约：red = sh + npos 而各行 npos 不同 ⇒ 启动时 smem 必须 ≥ dyn[1] + M + blockDim 个 float。M=1 时即现役口径。这条必须由 block e 执行（不许照抄 2*hd）。

证据一（内核级）：新增 tools/q35_attn_m.cu，对 m = 0..M−1 的「逐行」与「一次 M 行」两档比对：整个 KV 缓冲逐位相同、M 行输出逐位相同，外加一条独立正确性（追加进去的 M 行确实等于输入）。实测 16/16 全 PASS（M ∈ {1,2,4,8} × base ∈ {0,1,5,17}）。

假绿反证：把 q 的 M 行铺平（KV 新行保持互异）⇒ 各行的唯一差别就是「能看到多少 key」⇒ 行 0 与行 M−1 必须不同。若 npos 没跟着 m 走，各行会完全相同 ⇒ FAIL。

3.4.1 block c/d 的第二次翻车：自证自己退化了

首版把 KV 新行也铺平了，--base 0 当场报 FAIL。而那是断言退化、不是内核错：base=0 时前缀为空，各行看到的全是同一批相同向量 ⇒ softmax 均匀 ⇒ 各行都退化成 v，本来就该相同。

⇒ 纪律：写「必须不同」的自证之前，先想清楚差异是从哪来的、以及它在边界参数下还存不存在。

3.5 block e：加法式落地 + 一条会「砍掉自己验收能力」的绑定

改动面（刻意做成加法式）：

项做法
缓冲不碰任何现役缓冲；批量路径用自己的 slab d_mb（16 区域，行主序 [M][·]），惰性分配
入口新增 vcuda_dev_model_step_m / vllm_cuda_model_step_m；现役 step 一行未改
env 门VLLM_CUDA_Q35_MM=<M>（默认不设 ⇒ 入口直接拒）
/* 文件名：GitHupSRC/src/npu/cuda/vllm_cuda_kernels.cu（节选：vc_mb_ensure —— 惰性分配） */
static int vc_mb_ensure(vc_dev_t *s, int M) {
    if (!s || M < 2 || M > s->mdl_mmax) return -1;
    if (s->d_mb) return 0;                        /* 惰性：第一次真用 M 行入口才分配 */
    size_t per = 0;
    for (int i = 0; i < VC_MB_NREG; i++) per += vc_mb_width(s, i);
    if (per == 0) return -1;
    /* ★ 按 mdl_mmax（上限）分配，而不是本次的 M —— 行距是 mdl_mmax，必须容纳全部行。 */
    if (cudaMalloc((void **)&s->d_mb, per * (size_t)s->mdl_mmax * 4) != cudaSuccess) { ... }
    return 0;
}

端到端等价（本块的验收口径）：--mt M，A = 连续 M 次单 token、B = 一次 M 行：

M2345678
h_out / logits memcmp0000000

负向用例：--synthetic gdn --mt 4（当时 GDN 的 M 维还没做）⇒ 引擎拒绝。这正是设计原则的落点：没验证过的路宁可拒绝，不静默走。

★★ 过程中发现的一个真问题（会影响 A/B 方法本身）：连接器里那个 VC_BIND 是致命绑定——任一符号缺失就把整个 API 表判为失效。给 block e 加新入口时若按老办法写 VC_BIND，「新 exe 配旧 DLL」直接加载失败 ⇒ A/B 对照做不了了。

处置：新入口单取符号、取不到留 NULL，由 wrapper 拒绝：

/* 文件名：GitHupSRC/src/npu/cuda/vllm_cuda.c（节选） */
/* ★ block e：dev_model_step_m 是**可选**入口 —— **不走 VC_BIND**，
 *   因为它一旦缺失会把整个 API 表判为失效，而「**新 exe 配旧 DLL**」正是本项目
 *   A/B 对照的基础。取不到就保持 NULL，由 wrapper 判 NULL 拒绝。 */
*(void **)(&c->api.dev_model_step_m) = vc_lib_sym(c->lib, "vcuda_dev_model_step_m");

⇒ 记成规矩：凡新增的设备入口，都必须是「非致命绑定」，否则会悄悄砍掉自己的验收能力。（这条规矩后来在 MTP 身上又用了一次——第 31 篇。）

3.6 ② GDN 的 M 维：不做 chunked scan

递推在时间轴上必须串行，但权重只需读一遍（① 与 ⑦ 的投影都走批量）⇒ M 步只是串行 M 次状态更新：

段做法理由
① in_proj_qkv / in_proj_z批量同输入 ⇒ 权重只读一遍
②–⑥（conv+silu / qk_norm / delta / gate_norm）按时间串行 M 步，每步用第 m 行的指针递推在时间轴上必须串行；S/hist 是层内状态，串行就地推进 ⇒ 与「M 次单 token」是同一串操作
⑦ out_proj批量（M 行都准备好之后再发）权重读一遍
/* 文件名：GitHupSRC/src/npu/cuda/vllm_cuda_kernels.cu（节选：vc_gdn_step_dev_m） */
static int vc_gdn_step_dev_m(vc_dev_t *s, int layer, int M,
                             const float *d_h, float *d_out, cudaStream_t stream, ...) {
    ...
    /* ① 投影（**批量**：权重只读一遍） */
    if (vc_rq_launch_gemv_m(s, e_qkv, M, d_h, XC, stream) != 0) return -1;
    if (vc_rq_launch_gemv_m(s, e_z,   M, d_h, ZB, stream) != 0) return -1;
    /* ②–⑥ 按时间**串行 M 步**（`S` / `hist` 就地推进）*/
    for (int m = 0; m < M; m++) {
        const float *hm = d_h + (size_t)m * (size_t)dim;
        ...
        k_f32_gemv2<<<...>>>(w->d_in_a, w->d_in_b, hm, am, bm, vh, dim);
        k_gdn_beta_g<<<...>>>(am, bm, w->d_a_log, w->d_dt_bias, bem, gm, vh);
        k_gdn_conv_silu<<<...>>>(xcm, xcvm, hist, w->d_conv_w, conv_dim, ck);   /* hist 就地推进 */
        k_gdn_qk_norm<<<...>>>(...);
        vc_launch_gdn_delta(S, xcvm, xcvm + kw, xcvm + 2 * kw, gm, bem, onm,
                            kh, kd, vh, vd, stream);                            /* S 就地推进 */
        k_gdn_gate_norm<<<...>>>(...);
    }
    /* ⑦ out_proj（批量） */
}

端到端等价（三档配置全过）：

配置M=2M=3M=4M=6M=8
--synthetic gdn00000
--synthetic attn（回归）——0—0
★ 真实混合（48 GDN + 16 attn）0—0—0

★★ 第三行才是关键：--mt 现在能在真实混合配置上跑 ⇒ 「批量 verify 与逐行单 token 等价」已在真模型上端到端成立（不只是两种铺满的合成档）。

一处口径不对称（审计抓到）：预检用 vc_gdn_ready（允许 f32-only），而步进要求三个 rq 键齐备 ⇒ f32-only 配置能过预检、却在 body 第 N 层失败，而前面各层的 KV 已经写进去了。⇒ 已把「三键齐备」并入预检，与步进同口径。

3.7 ④ 状态克隆 / 回滚：两种状态，两种退法

投机要「试跑 M 步后回滚」。两个入口 mark（快照）/ rollback（恢复），只搬状态、不发计算内核：

状态怎么回滚为什么
GDN S（149.6 MB）与 histD2D 快照 / 恢复它们是设备上就地覆写的，没有别的办法
attn 的 KV 行不拷，只把宿主游标改回去可见长度 npos = dyn[1]+m+1 由游标派生 ⇒ 多写的行永远读不到，下次 append 会覆盖它 ⇒ 零拷贝
/* 文件名：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));  /* ★ attn KV 只退游标、零拷贝 */
    if (s->mtp_len_snap >= 0) { s->mtp_len = s->mtp_len_snap; s->mtp_trace_done = 0; }
    s->mdl_mk_ok = 0;      /* ★ 一次性：回滚后快照失效，必须重新 mark 才能再回滚 */

那条「attn 的 KV 不用拷」的论证，与第 19/20 篇 KV 镜像的「游标即可见长度」是同一条。换成投机场景，它变成了「回滚免费」的理由。

★ 快照是一次性的；★ 门控与批量路径同源（要求 MM ≥ 2）；快照惰性分配（不用就不占显存）。

证据（有区分度，不是自证）：--rb M

  • A 档：reset → 连续 M 次单 token
  • B 档：reset → mark → 一次 M 行（把状态占掉）→ rollback → 再跑 M 次单 token
  • 判据：ref == aft 逐字节
配置M=2M=4M=8
★ 真实混合000
--synthetic gdn000

⇒ 若回滚没生效，B 档第二次单 token 会读到被推进过的 S/hist，结果必然与 A 档不同——这是一个真正有区分度的判据（不是「回滚后重跑还是原来那批输出」那种自证）。

mermaid diagram

图 1：M 行批量前向的五个块（a–e）与「退回去」的两个原语。五个块依次是行内 norm/add/swiglu、按 R 表投影批量、RoPE 行 m 用位置 pos+m、KV append + attention（行 m 只看 prefix+m+1 个 key）、以及 env 门默认关的批量入口；两种层型的批量难点分别落在 GDN（时间轴串行 M 步）与 full-attn（块内因果）。mark/rollback 只搬状态、不发计算内核，且两种退法不同：GDN 的 S/hist 就地覆写只能 D2D 快照/恢复，attn 的 KV 行可见长度由游标派生 ⇒ 退游标、零拷贝。全部判据 memcmp = 0。

四、实测数据

4.1 尺寸轨迹与「护栏」的演进

尺寸轨迹：block b 1,815,040 → c 1,817,088 → c+d 1,818,112 → c+e 2,472,960 → c+d+e 2,473,472 → e+gdn 2,475,520。

block e 那一跳很大，是因为 10 个 k_rq_q4_gemv_mr<R,M> 模板实例化。附带一个诚实标注：热路径 +0.33% 的三轮同号，但 n=3 不足以判显著（符号检验 p≈0.125）⇒ 按纪律暂记在噪声带内，不列为已确认的回归。

护栏从「纪律」升级成「硬失败 + 不可变归档」：block c/d 那一轮忘了刷新 build-x64-prev\ ⇒ 对照实际是「block b vs block c+d」的累计，丢掉了单块增量证据。审计指出「别只写成纪律」，于是补三条：

#护栏形式
①′旁路目录是「上一次已验证状态」的滚动副本；每落一块之前先刷新纪律
①″跑 A/B 之前先校验旁路 DLL 的 SHA == 上次验收记录的 SHA；不等就拒绝跑硬失败
①‴另存不可变归档（只增不改），避免「预置基线被自己覆盖」归档

4.2 汇总读数

判据口径结果
--mt M 端到端等价M ∈ {2..8}h_out/logits 全 memcmp=0（逐位）
内核级等价M ∈ {1,2,4,8} × base ∈ {0,1,5,17}16/16 PASS
--rb M 回滚等价真实混合 / gdn，M ∈ {2,4,8}全 memcmp=0
负向用例--synthetic gdn --mt 4（GDN M 维未做时）引擎拒绝

五、边界与已知限制（必须原样保留）

  1. 覆盖：只有 --synthetic attn / gdn / 真实混合三种配置（后来含 GDN 也通了）。
  2. 只走流式；graph 档直接拒绝（grid/smem 在捕获时定死）。
  3. 批量档只走 rq 池；f32 投影拒绝；不做 qkv3 融合（那是 M=1 的 #0③）。
  4. ★ 已证的是「前向输出逐位等价」（h_out/logits）；步后状态（S/hist/attn_len）与续跑等价未验证（测试两端各 reset 一次，没验连续性）。
  5. 收益门禁：M=1 三文件全等 ✅；端到端收益要等接线（⏳）。⇒ 本块不产出任何收益结论，也不能说「投机已经能跑」——它证明的是**「批量前向的算术是对的」**，没证明「投机更快」。

CPU 对照（迁移前基线）

  • CPU 参考：model/vllm_q35_model.c:q35_model_step / q35_model_prefill（整步编排）、model/vllm_q35_gdn.c:q35_gdn_step（GDN 递推）、model/vllm_q35_attn.c:q35_attn_step（全注意力）—— 逐行单 token 的 f32 参考语义。
  • 迁移要点：新建 M 行批量前向（投影 + norm + GDN + full-attn），要求与「逐行单 token」逐位等价；M 由 gridDim.x 派生（不加内核参数，避开 graph 捕获定死），attn 引入块内因果 npos=dyn[1]+m+1，状态用 GDN 快照 / attn 退游标两种退法。
  • 真机验证：命中 E4、E5；以及 E2 的 M 行批量路径（前向 14.28 ms/token、位级 logits maxae=6.1e-05 maxrel=3.08e-06 OK、贪心 64/64）。内核级 tools/q35_attn_m.cu 16/16 与 --rb 回滚等价未独立复现：属工具档（需 VLLM_CUDA_Q35_MM 门），不在本机清单。

六、小结

  1. 「加法式落地」让「不启用就不变」变成结构性质。不碰现役缓冲、现役入口一行未改、默认档连一次分配都不多 ⇒ 不需要靠人工核对去保证零回归。这是第 26/27/29 篇「默认关 + 逃生舱」在工程结构上的延伸。
  2. 两种状态、两种退法，是「有状态算子」在投机场景下的核心区分：attn 的 KV 状态是可寻址的（游标即长度）⇒ 退游标、零拷贝；GDN 的递推 S 状态是就地叠出来的 ⇒ 只能快照/恢复。⇒ 第 32 篇那道「退不回去的墙」，根就在这里。
  3. 判据要「有区分度」，不是「能通过」。--rb 之所以值钱，是因为它若回滚失效就必然变红。反过来，block d 那条「必须不同」的自证在边界参数下退化了——说明自证的效力也是有限的，要连它的失效条件一起想。
  4. 绑定的宽严是工程决策，不是风格问题。VC_BIND 的「一个都不能少」对主链路是对的，对后加的可选入口是错的——写错了会砍掉自己的验收能力（新 exe 配旧 DLL 直接加载失败，连「新旧对照」都做不了）。⇒ 可选入口一律非致命绑定。
  5. 基线必须留档，且要能防自己。「忘了刷新」这种失误不能靠纪律：纪律会被忘记，硬失败断言 + 不可变归档不会。

相关篇目：第 32 篇（用这些基建跑投机，以及为什么判否）；第 31 篇（复用同一套「非致命绑定」规矩）；第 19 篇 / 第 20 篇（「游标即可见长度」的来源）；第 35 篇（在这条批量路径上继续「把 m 循环搬进内核」）。 形状门禁 / R 行 / 五块工单与尺寸轨迹的细则见仓库 GitHupSRC/docs/。 源码与配套资源：本仓库 gitee.com/pei-xiaogua…； CPU 推理源码 gitee.com/pei-xiaogua…

avatar
文章
阅读
粉丝