50 · prefill 追赶开篇：把 15 倍的账拆开

一句话：15.4× 不是一个能放着不管的数；本篇重新开账 —— 不做内层微优化、也不接受「混合模型 ⇒ GDN 是瓶颈」的直觉，而是先把 15.4× 拆成几个桶：实测否掉了「限速项是指令发射」，「首要环节是投影 GEMM（相对差 18.5×），不是 GDN」，于是候选收成三个（A 分块降重读 / B int8 / C 张量核）。
前置：建议先读 第 47 篇（15.4× 的来源）与 第 48 篇（廉价路径封顶 —— 天花板 ≈4.7×）。
环境：x86-64 + NVIDIA RTX 5090 · Qwen3.5-27B（.vqf q4）· 探针口径为 L2 档（权重 <96 MB）。

一、问题与结论

第 47 篇把 llama.cpp 请到同一块卡上，拿到本轮最不好看的一个数：prefill 512：本引擎 ≈274 t/s vs llama.cpp 4206 t/s ⇒ ≈15.4×（那一篇只报数、不归因）。第 48 / 49 篇把「改个内核」的直觉逐条量死，并算出廉价路径的天花板 ≈4.7×（远小于 15.4×），于是这条线在微优化上收口为「不做」。

可 15.4× 不是一个能放着不管的数。这一篇是重新开账：不做内层微优化、也不接受「混合模型 ⇒ GDN 是瓶颈」的直觉，而是先把 15.4× 拆成几个桶，看每个桶各差多少、哪个桶值得动。

结论	数值	标注
限速项不是「指令发射」	dp4a 把每 MAC 指令数砍掉数倍却更慢；权重带宽只到 HBM 的 19–26%	实测
dp4a 出局	与显式 FMA 版同速，却多付激活 int8 量化、精度 RMSrel 差 3–4 个数量级	实测
跨引擎分桶	GEMM ≈18.5×、GDN ≈11.6×、总 ≈15.3×	实测（nsys）+ 归一化
首个环节	投影 GEMM（相对差 18.5×），不是 GDN	实测

逐条结论：

1. 限速项不是「指令发射」（Phase 0 探针小样）：把每 MAC 的指令数砍掉数倍的 dp4a 变体跑得更慢；实测权重带宽只到 HBM 的 19–26%。⇒ 方向从「减指令」改成「降延迟 / 提占用率」。
2. dp4a 出局（同一轮）：与显式 FMA 版同速，却要多付激活 int8 量化，精度 RMSrel 差 3–4 个数量级。
3. 跨引擎分桶把账拆开了：GEMM ≈18.5×、GDN ≈11.6×、总 ≈15.3×。⇒ 先打 GEMM；「GDN 是瓶颈」的直觉被实测否掉（它还是第二位）。
4. llama.cpp 那边的机制读出来了（源码/内核名证据）：一趟算完 512（我方 ≈22× 的发射/重读次数）、int8 MMQ（我方 f32 + FP32-FMA，只到 FP32 峰值的 ≈18%）、stream-K（我方无）。
5. 于是候选收成三个：A 分块降重读 / B int8 / C 张量核。用户裁定先走 A。★ 结果 A 在下一篇就被实测否掉，B 在再下一篇否掉 —— 这条线的真正出口是 C。

二、Phase 0：先花最小成本，把「限速项」定下来

按项目纪律，动引擎之前先做探针小样（引擎零改动）。给 tools/q35_gemv_mb.cu 加两个非位级变体。

★ 探针自身的缺陷先记一笔：首跑 dp4a 出 maxrel=0.00 / nan。根因是探针原先用 fill_xor 填随机字节当权重，把 72 B 瓦片里的 f16 标度填成了 NaN/Inf（约 6%）⇒ 精度对照整片失真。补 fill_q4（标度取小正数）后重跑，本节全部数字出自修后的日志：

/* tools/q35_gemv_mb.cu:103-137 —— 有效 q4 权重填充（修掉 NaN/Inf 标度） */
static void fill_xor(uint8_t *p, size_t n, uint32_t s) {
    for (size_t i = 0; i < n; i++) { s ^= s << 13; s ^= s >> 17; s ^= s << 5; p[i] = (uint8_t)s; }
}

/* ★ Phase 0：有效 q4 权重填充（V4/V5 要出 maxrel ⇒ 随机字节不行）。
 *   原 `fill_xor` 会把 f16 标度的位模式填成 NaN/Inf（~6%）⇒ 精度对照整片失真
 *   （V5 曾出现 maxrel=0.00/nan）。本函数按现役 72 B 瓦片布局填：
 *   d[4] = 4 个小正数 f16（0.002–0.020），其余 64 B 为随机 nibble。 */
static void fill_q4(uint8_t *p, size_t nbytes, uint32_t s) {
    for (size_t t = 0; t + VC_RQ_TILE_BYTES <= nbytes; t += VC_RQ_TILE_BYTES) {
        for (int r = 0; r < 4; r++) {
            s ^= s << 13; s ^= s >> 17; s ^= s << 5;
            const float v = 0.002f + 0.018f * (float)((s >> 8) & 0xFFu) / 255.0f;
            const uint16_t h = f32_to_f16(v);
            p[t + r * 2 + 0] = (uint8_t)(h & 0xFFu);
            p[t + r * 2 + 1] = (uint8_t)(h >> 8);
        }
        /* 其余 64 B：随机 nibble */
    }
}

V4 = 显式 FMA 收缩：权重解包外提后内层用 fmaf() ⇒ 每条 MAC 1 条 FFMA，而基线在 -fmad=false 下是 FMUL + FADD 两条：

/* tools/q35_gemv_mb.cu:347-405 —— V4：显式 FMA 变体（解包外提 + fmaf） */
template <int R, int M>
__global__ void k_mb_fma(const uint8_t *__restrict W, const float *__restrict x,
                         float *__restrict y, int rows, int cols) {
    ...
    for (; b + 8 <= nb; b += 8) {
        float wlo[R][4], whi[R][4];              /* ★ 权重解包**外提**、常驻 */
        for (int r = 0; r < R; r++) { /* 解包 nibble → wlo/whi（含 s 标度） */ }
        for (int m = 0; m < M; m++) {
            /* 取 x 的 8 个分量 xs[0..7] */
            for (int r = 0; r < R; r++) {
                float t0 = 0.0f, t1 = 0.0f;
                for (int k = 0; k < 4; k++) {
                    t0 = fmaf(wlo[r][k], xs[k],     t0);   /* ★ 每条 MAC 只 1 条 FFMA */
                    t1 = fmaf(whi[r][k], xs[4 + k], t1);
                }
                acc[r][m] += t0 + t1;
            }
        }
    }
}

V5 = dp4a：nibble→int8 用 3 条无跨字节借位的整数指令；激活每 32 列块对称量化到 int8（计时区外准备）：

/* tools/q35_gemv_mb.cu:434-492 —— V5：dp4a 整数点积（每 (r,m) 每块 2 条 dp4a） */
/* 前提（计时区外）：激活 xq 每 (m, 32 列块) 对称量化到 int8，dx[m][nb] 为对应 f32 标度；
 *   权重保持现役 q4 打包（nibble 解包仍在计时区内，如实计入）。
 *   内层 nibble(0..15) → 有符号 int8，用 3 条整数指令且无跨字节借位：
 *     i8 = ((nib ^ 0x08080808) + 0x78787878) ^ 0x80808080
 *   ⇒ 每 (r,m) 每块 2 条 dp4a 替代 8 次 (I2F + FMUL + FADD)。 */
...
    int sumi = __dp4a(wlo[r], xlo, 0);
    sumi     = __dp4a(whi[r], xhi, sumi);

结果（M=16，单位 ms）：

形状	权重 MB	基线 R=2	V4 R=4	V5 R=4	最好/基线	标注
gate/up 17408×5120	47.8	0.150	0.104	0.107	1.44×	实测
qkv 10240×5120	28.1	0.090	0.053	0.064	1.70×	实测
q 12288×5120	33.8	0.114	0.092	0.087	1.24×	实测
down 5120×17408	47.8	0.138	0.133	0.234	1.04×	实测
o 5120×6144	16.9	0.052	0.061	0.084	0.85×（更慢）	实测

精度（相对基线）：V4 元素 maxrel 5.9e-2（大投影）～3.5e-4；V5 的 maxrel 高达 4.0e2（meanrel ≈ 3e-2）。

判定：

• P1（显式 FMA）= 正。价值不在 R=2（0.85–1.24×，符号不一），而在解锁 R=4 @ M=16 —— 大投影 1.44–1.70×。
• P2（dp4a）= 负且被支配：R=2 档普遍更慢；R=4 档与 V4 持平，却要多付激活量化、精度差 3–4 个数量级。

★ 本轮最重要的负结果：dp4a 把指令/MAC 砍了数倍却更慢 ⇒ 这个内核不是「指令发射受限」。实测权重带宽只到 335（基线）→ 460（V4）GB/s，即 HBM 峰值的 19–26% ⇒ 真正的形态是「每 token 的工作量导致占用率 / 延迟隐藏不足」，指令数只是它的一个代理量。

★ 这一条顺手订正了第 48 篇的表述（那里把限速项写成「每 token 的指令发射」）。也订正了「M=16 时 R=4 一律慢 1.7–2.0×」这条边界 —— 它只对无 fmaf 的基线成立：V4 下 R=4@M=16 反而最快。⇒ 那不是寄存器的硬墙，而是与内层指令形态耦合的。

★ 边界：探针只出 GEMM 时间 + 元素级 maxrel；未做引擎侧 A2/A3 ⇒ 尚不能判 P1 能否过精度闸门。且元素级 maxrel 会被「点积抵消」放大（V4 仅改结合序就有 5.9e-2）⇒ 精度阈值不能直接取元素 maxrel，须另定口径。

三、跨引擎分桶：把 15.4× 拆成「谁占了多少」

Phase 0 之后的第一件事不是动手，而是把账拆开。两侧各自跑 nsys，比桶的方向与量级。

★ 口径前提：两侧量化格不同（llama.cpp Q4_0 有效 4.60 bpw vs 本引擎 .vqf q4 ≈4.25 bpw），激活精度也不同 ⇒ 本节只比桶的方向 / 量级，不外推绝对倍数、不做桶级绝对归因。

基线先复现：llama.cpp pp512 = 4167.14 ± 234.67 t/s（复现第 47 篇的 4206 ± 220）；本引擎 --perf-partA --bench-seqlen 512（门关）= 1.876 s ⇒ 273 t/s（复现第 47 篇的 274）。

各自的分桶（nsys cuda_gpu_kern_sum）：

桶	本引擎（kernel 总 9093 ms†）	占比	llama.cpp（kernel 总 113.9 ms）	占比	标注
GEMM	7639 ms（rq_q4_gemv_mr<2,16> 3707 + mr3<2,16> 3921）	84.0%	79.0 ms（mul_mat_q + stream_k_fixup）	69.3%	实测（nsys）
GDN	1119 ms	12.3%	18.6 ms	16.3%	实测（nsys）
attn	244 ms	2.7%	0.8 ms	0.7%	实测（nsys）
norm/act	211 ms	2.3%	8.1 ms	7.1%	实测（nsys）

† 我方 nsys 覆盖 6 趟 prefill + 6×40 decode ⇒ 需归一化。

归一化到「一趟 prefill 512」后：

桶	本引擎	llama.cpp	相对差（我方/对侧）	标注
GEMM	≈1576 ms	≈85 ms	≈18.5×	实测 + 归一化
GDN	≈231 ms	≈20 ms	≈11.6×	实测 + 归一化
总	1876 ms	123 ms	≈15.3×	实测 + 归一化

★ 正确读法：不是「GDN 是瓶颈」（它还是第二位），也不是「再做一次内层微优化」—— 是**「把 GEMV 骨架换成真 batched / 分块 GEMM」**。

llama.cpp 那边是怎么做的（内核名证据，不是推断）：

观察	llama.cpp	本引擎	标注
一趟算完 512	mul_mat_q<...> 400 次/趟（≈ 每投影一次）	mr<2,16> + mr3<2,16> ≈ 9000 次/趟 ⇒ ≈22×	实测（nsys）
计算路径	quantize_mmq_q8_1 496 次 ⇒ int8 MMQ	f32 激活 + FP32-FMA（≈18% FP32 峰值）	实测（nsys）
分块	mul_mat_q_stream_k_fixup 320 次 ⇒ stream-K	无	实测（nsys）

★ 两个结构成因：① 每趟 token 数 M（16 vs 512） ⇒ 权重重读 ≈32×；② 计算路径（f32/FP32-FMA vs int8-MMQ）。

图 1：15.4× 的账怎么拆开。从左到右读：起点是「prefill 512 = 274 vs 4206 t/s（≈15.4×）」；Phase 0 探针先否掉「指令发射」这个限速项（dp4a 指令更少却更慢）；跨引擎分桶（归一化到一趟）得到总 ≈15.3×，其中投影 GEMM ≈18.5×（占 84%，首个环节）、GDN ≈11.6×（占 12.3%，第二位），attn（2.7%）与 norm（2.3%）是小桶；据此候选收成 A/B/C，用户裁定先走 A。

四、候选与裁定

候选	直击	位级	备注
A. 真 batched GEMM（smem 分块 / 大 M）	成因 1（+2 的一部分）	可能位级可保（M/R 不改行内求和序）	换访问模式；改动最大但直击 4.7× 上界
B. int8 路径（大 M 下重测）	成因 2	非位级	dp4a 的负结果只在 M=16 测过；大 M 下结论可能不同
C. 张量核	成因 2（换天花板）	非位级	需改布局

用户裁定走 A。本篇到此收口 —— 但 A 的结局在下一篇，而且不是好消息。

五、边界与已知限制

• 权重全 L2 档（<96 MB），不是整模型冷档 ⇒ 探针只作相对对照。
• nsys 的 Time (%) 含 profiler 扰动；两侧口径不同 ⇒ 只比方向与量级。
• 本轮只动 tools 探针 + 文档（引擎 src/** 与 DLL 未动），未改任何默认值。

CPU 对照（迁移前基线）

• CPU 参考：src/model/vllm_q35_fused.c:q35_fused_gemv —— 融合量化 GEMV 的 CPU f32 参考，探针与对拍的同语义基线。
• 迁移要点：跨引擎 nsys 分桶把 prefill 15.4× 拆开，定位首个环节是投影 GEMM（≈18.5×）而非 GDN；dp4a 负结果把限速项从「指令发射」改判为「占用率／延迟隐藏」。
• 真机验证：未独立复现：探针为 L2 档（权重 <96 MB）单读数；两侧量化格不同（llama.cpp Q4_0 4.60 bpw vs .vqf q4 ≈4.25 bpw），只比方向与量级。

六、小结（可复用结论）

1. ★ 动手前先把「整条差距」拆成桶，再决定先打哪一个。 分桶把「GDN 是瓶颈」的直觉直接推翻了 —— 首个环节是投影 GEMM（相对差 18.5×），而 GDN 只是第二位。
2. 负结果要量的更细一点：dp4a「指令更少却更慢」，把限速项从「指令发射」改判为「占用率 / 延迟隐藏」—— 这是后面几轮所有优化的方向依据。
3. 探针自身的缺陷要当场修、当场重采。权重填充把 f16 标度填成 NaN/Inf，会整片污染精度对照；补 fill_q4 后本节数字才成立。

---

相关篇目：第 47 篇（把 llama.cpp 请上卡 —— 15.4× 的来源）、第 48 篇（廉价路径封顶 —— 天花板 4.7×）、第 49 篇（投机解码拆账 —— 同一套「先拆账再动手」的方法）。仓库 docs/ 的适配规格中有对应的 Phase 0 探针小样与跨引擎分桶记录。
源码与配套资源：本仓库 https://gitee.com/pei-xiaoguang/kestrel-llm-cuda.git；
CPU 推理源码 https://gitee.com/pei-xiaoguang/kestrel-llm

下一篇：51 · 两条廉价路径逐条关闭：占用率与 dp4a