Kestrel_VLLM - 40 · 把 CUDA Graph 默认打开：一条默认值改动的两个前置

40 · 把 CUDA Graph 默认打开：一条默认值改动的两个前置

一句话：hybrid 的 CUDA Graph 早已落地、位级也证过，但它一直默认关 ⇒ 生产上从没开过。把默认值从关改成开看起来是一行的事，实际是「先量 → 拆门 → 默认化 → 全回归」：其中一个自己的推断被先量推翻，还有一道门必须先拆，不拆就会在第一次跑回归时把 –mt 当场打断。 前置：第 29 篇（CUDA Graph 接到 hybrid —— 图化的真正工作在「把随 token 变化的量赶出内核参数」）、第 36 篇（一次退役 —— M 行 graph 为什么停用）。 环境：x86-64 + NVIDIA RTX 5090 · Qwen3.5-27B q4（hybrid）· CUDA 后端（阶段五）。

一、问题与结论

hybrid 的 CUDA Graph 早就落地了，而且位级已经证过（Graph 档 vs 流式档的 h64/hfinal/logits 逐字节相同）。但它一直默认关 ⇒ 生产上从没开过。

「把默认值从关改成开」看起来是一行的事。实际不然：这一轮里，有一个我自己的推断被先量推翻，还有一道门必须先拆 —— 不拆，默认化会在第一次跑回归时把 –mt 当场打断。

要点：

二、背景

graph 的门一旦默认开，最先被问的就是「短序列会不会更亏」——图捕获有固定成本，序列越短，这笔摊得越薄。所以在改默认值之前，先用一个零改动探针把这个问题量掉；量完才发现，真正要先处理的是这个门还管着什么别的。

三、核心机制与做法

3.1 先量：短 T 到底亏不亏

q35_gpu_model –iters N（A = graph 关；B = graph 开；同轮内配对）：

⇒ ≈1.11×，且与 N 无关。N=5 也是 1.11× ⇒ 「短 T 净亏」不成立，阈值门基本不需要。

★ 自纠（这一条是本轮最重要的方法论点）：我此前据「引擎 forward 分支在 T=5 时 ON 反而慢 ≈10.1 ms」推断要加「按 T 的阈值门」。复测才看清：

那不是 graph 变慢，而是捕获成本被算进了该路径的被测窗口（prepare 没有 reset）。一旦捕获落在请求窗口之外（比如以 reset 为序列起点），这是一次性的。

这条后来被写进判据 G3：捕获必须在请求窗口外。

3.2 为什么默认化之前必须先拆门

现状：一个门 VLLM_CUDA_Q35_GRAPH 同时管两件事 —— 「M=1 步进图」与 M 行 per-M 图（代码注释：「沿用同一个门」）。

而 vcuda_dev_model_step_m（M 行入口）在 s->mdl_graph_mode 为真时直接 return -1（批量路径只走流式）：

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu —— M 行入口：graph 模式下直接拒绝 */
VCUDA_API int vcuda_dev_model_step_m(void *dev, int M, int pos,
const float *h, float *h_out, float *logits) {
vc_dev_t *s = (vc_dev_t *)dev;
if (!s || !s->ok || !h) return –1;
if (M < 2 || M > s->mdl_mmax) return –1; /* 门没开 / M 不匹配 ⇒ 拒绝 */
/* … */
if (s->mdl_graph_m_mode) return –1; /* 批量路径只走流式 */
/* … */
}

⇒ 把默认值一改成「开」，M 行入口默认就拒 —— –mt 与全部 M 行回归当场失效。这不是「可能出问题」，是必然。

做法（拆门）：把两个特性拆成两个独立门，各读各的 env：

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu —— 拆门后的两个 env 读取 */
/* M=1 的图：§18.40 起**默认开**，只有显式 `0` 才关。 */
static int vc_mdl_graph_env(void) {
static int v = –1;
if (v < 0) {
const char *e = getenv("VLLM_CUDA_Q35_GRAPH");
v = (e && e[0] == '0') ? 0 : 1; /* 默认开，=0 才关 */
}
return v;
}
/* M 行 per-M 图（已退役）：**独立门**，**默认关**，只有显式开才开。 */
static int vc_mdl_graph_m_env(void) {
static int v = –1;
if (v < 0) {
const char *e = getenv("VLLM_CUDA_Q35_GRAPH_M");
v = (e && e[0] && e[0] != '0') ? 1 : 0; /* 只有显式开才开 */
}
return v;
}

（上面代码块里的 §18.40 是仓库内部规格编号，读者可忽略；详见仓库 docs/。）

vcuda_dev_model_step_m 的拒绝改判新门（新字段 mdl_graph_m_mode）⇒ 新门默认关 ⇒ –mt 行为不变。

这里的分工很干净：默认化只动 M=1 那条；M 行图（已退役）只把它的门独立出来，不重启它。

3.3 判据（预登记，写死，防事后改判）

3.4 审计抓到一个拆门漏项

拆门不是「把拒绝判据改一下」就完了 —— 任何以 mdl_graph_mode 为 latch 的地方都要重新判。

漏掉的是 vc_attn_step_dev_m：它按 s->mdl_graph_mode 决定动态 smem 的定尺与 cudaFuncSetAttribute 守卫。若不改，M 行路径会因 M=1 图默认开而静默改用 graph 档定尺，同时把属性守卫变成死代码。这两处的 latch 必须改判 M 行自己的门：

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu —— 现役内核的动态 smem 定尺（改判 M 行门） */
if (!fa_on) { /* 现役内核：门关 / FA 未启用 / 被门限跳过 */
const size_t shbytes = s->mdl_graph_m_mode
? ((size_t)s->attn_maxseq + (size_t)M + (size_t)hd) * 4 /* graph 档：与 pos 无关 */
: ((size_t)(len + M) + (size_t)hd) * 4; /* 流式档：按当前 len */
if (!s->mdl_graph_m_mode && shbytes > 48 * 1024) {
if (cudaFuncSetAttribute(k_attn_decode,
cudaFuncAttributeMaxDynamicSharedMemorySize,
(int)shbytes) != cudaSuccess) {
cudaGetLastError(); return –1;
}
}
k_attn_decode<<<M * nh, hd, shbytes, stream>>>(mao, mbg, kc, vc, s->d_dyn, /* … */);
}

订正：这两处的 latch 由 mdl_graph_mode 改判 mdl_graph_m_mode。

★ 附带一个必须记下的前置条件：vc_attn_step_dev_m 里 mdl_graph_m_mode ≡ 0 是载荷不变量，仅由「vcuda_dev_model_step_m 开头 early-return」+「vc_mdl_graph_m_ensure 不可达」共同守护。若将来有人新增一个「开新门却绕过 ensure」的 vc_model_step_body_m 直调者，该处会走 graph 档定尺却跳过 cudaFuncSetAttribute ⇒ shbytes > 48 KB 时可能发射失败。重接那条 M 行 graph 之前必须先补这条。

同轮还顺带订正了 3 处过期注释（两处「默认关」→「默认开」；G2 比值注释补上「终档 38.6×」），保持行数不变，故台账行号不受影响。

四、实测数据

实测（最终档）：

拆门的端到端回归（证明「–mt 行为不变」）：默认档（graph 开）跑 VLLM_CUDA_Q35_MM=4 –mt 4 ⇒ h_out memcmp=0、logits memcmp=0、等价 PASS、M 行提速 2.325×（与拆门前一致；未被默认开牵连拒绝）。

全回归 8/8（默认化触及所有 M=1 步进路径，按纪律全跑）：

⇒ 8/8 逐项与 _w4g 相同。

五、边界与已知限制

• 收益 ≈1.11× 是同一 harness 内的同会话交错 A/B；绝对值随 GPU 时钟/同会话抖动变化，只作方向判断（引用时必须带日志名）。
• G3 的「零劣化」只对以 reset 为序列起点的调用方成立；引擎 forward/prepare 那条仍在一次捕获成本（≈10 ms）以内，不是零劣化 —— 口径必须这样写。
• M 行 graph 本轮不重启：只把它的门独立出来。
• 与 CPU 的位级一致：有意不做（沿用阶段五口径纪律）。

CPU 对照（迁移前基线）

• CPU 参考：CPU 侧无对应实现（CUDA Graph 属 CUDA 后端同期新增）—— CPU 侧即逐 token 步进。
• 迁移要点：把「一个门管两特性」拆成 VLLM_CUDA_Q35_GRAPH（M=1，默认开）与 VLLM_CUDA_Q35_GRAPH_M（M 行 per-M 图，默认关）；所有以旧 latch 为准处逐一改判；捕获须在请求窗口外。
• 真机验证：命中文中 8/8 回归项中的 E7（–test-l3 PASS(0)、–test-sparse PASS(9)）；graph 默认开后 E2 于默认档运行。≈1.11× 收益为同会话交错 A/B 读数，未在本机清单复现。

六、小结

相关篇目：第 29 篇（CUDA Graph 接到 hybrid）；第 36 篇（一次退役 —— M 行 graph 为什么停用）；第 41 篇（prefill 批量化 —— 承接本条，把 M 行路径接到 prefill）。 源码与配套资源：本仓库 https://gitee.com/pei-xiaoguang/kestrel-llm-cuda.git； CPU 推理源码 https://gitee.com/pei-xiaoguang/kestrel-llm