[Bug][DCP] GLM-5.3 dense prefill consumes uninitialized K rows on non-owner ranks

快速结论：该报错发生在 vLLM 使用 DCP（解码上下文并行）运行 GLM-5.3 等模型的 dense prefill 阶段，非 owner rank 因 fused_norm_rope 内核的 slot_mapping 负值早退条件错误，导致 K 矩阵（kv_c_out、k_pe_out）未初始化就被 dense prefill 消费。优先排查 models/deepseek_v32/common/kernels.py 中 fused_norm_rope 内核的负 slot 早退逻辑是否缺少“是否正在物化 K 输出”的额外判断。

适用环境：vLLM（Issue 中验证版本为包含 #50005 修复的 cu130 构建及 current main ee3c00bb），NVIDIA GB200（4x），GLM-5.3-NVFP4 模型，TP4/DCP4，FP8 KV cache，启用了 chunked prefill、prefix caching 和 CUDA graphs。操作系统、Python、CUDA 具体版本未在 Issue 中明确列出。

最快修复方案：将 Issue #54908 中的补丁应用到 vllm/models/deepseek_v32/common/kernels.py（对应部署镜像为 vllm/vllm-openai:nightly-7c5dc571cbd1064ecc8a9b1045637ff647aa22cb），核心是修改 fused_norm_rope 内核中 slot_mapping 为负时的早退条件，使 K 输出在每个 rank 上都正常物化，同时保留对 owner 本地缓存写操作的负 slot 守卫。若当前环境已包含 #54908 修复，则无需额外操作。

注意事项：该补丁在 Issue 验证中通过 CUDA graph 矩阵和 67/67 内核测试，但补丁发布时仍标记为“human submitter will review”，建议先在非生产环境验证；评论者部署后“I’ll get feedback on whether it breaks or not”，说明尚无大规模生产环境的长期稳定性验证。

问题场景

用户在使用 vLLM 部署 GLM-5.3-NVFP4 模型时，采用 TP4/DCP4（tensor parallel 4 + decode context parallel 4）配置，启用 FP8 KV cache、chunked prefill、prefix caching 和 CUDA graphs。在 dense prefill 阶段，模型输出的 token 与 TP4/DCP1（仅 tensor parallel 4）配置相比出现确定性的输出不一致（deterministically incoherent），即使应用了此前的 #50005 正确性修复并禁用了直接 DCP 通信路径，问题仍然存在。

报错原文

[Bug][DCP] GLM-5.3 dense prefill consumes uninitialized K rows on non-owner ranks
fused_norm_rope kernel currently returns when slot_mapping is negative. Under DCP, a negative local slot can mean that another rank owns the cache row; it does not mean that the dense-prefill K row is unused. As a result, non-owner ranks leave kv_c_out and k_pe_out uninitialized, and dense prefill consumes those buffers.

原因分析

可能原因是 fused_norm_rope Triton 内核的早退条件设计存在缺陷。在单 rank（如 TP4/DCP1）环境下，slot_mapping 为负值确实意味着当前 token 的 cache 行未被使用，此时可以直接返回。但在 DCP（decode context parallel）环境下，每个 rank 只管理部分 cache 行，非 owner rank 的 slot_mapping 可能为负，但这仅表示 cache 行由其他 rank 拥有，并不意味着 dense prefill 中需要物化的 K 行不存在。NVIDIA DeepSeek-V3.2 / GLM-5.3 的 dense prefill 需要在所有 DCP rank 上执行归一化/旋转后的 K 物化（normalized/rotated K materialization），而当前内核在 slot_mapping 为负时直接返回，导致 kv_c_out 和 k_pe_out 缓冲区未写入任何数据，后续 dense prefill 读取这些未初始化缓冲区产生错误结果。

核心在于 cache 所有权（cache ownership）与 K 物化（K materialization）需要两个独立的守卫条件：K 输出需要在每个 rank 上物化，而 owner 本地 cache 写入前的负 slot 守卫需要保留。当前代码把这两个场景混为一谈，导致非 owner rank 跳过 K 物化。

环境排查

确认 vLLM 版本是否包含 #50005 修复，以及是否在当前 main 分支或接近版本（Issue 中当前 main 为 ee3c00bbf47e0ef7e975705cc980b06ee5576bb0，部署验证镜像为 vllm/vllm-openai:nightly-7c5dc571cbd1064ecc8a9b1045637ff647aa22cb）。
确认是否启用了 DCP：--decode-context-parallel-size 大于 1（如 4），此时才会出现该问题。
确认是否运行 GLM-5.3 或 DeepSeek-V3.2 系列模型并执行 dense prefill（如长上下文、无 cache 命中场景）。
检查 vllm/models/deepseek_v32/common/kernels.py 中 fused_norm_rope 内核的早退条件代码是否为旧逻辑。

解决步骤

定位 vllm/models/deepseek_v32/common/kernels.py 中 fused_norm_rope 内核代码。
应用 Issue #54908 补丁，修改两处逻辑：
- Hunk 1：在 slot_mapping 为负的早退条件中增加额外判断：elif tl.load(slot_mapping_ptr + tok_idx) < 0 and (pid != 1 or (kv_out_ptr is None and kpe_out_ptr is None)):，即在需要物化 K 输出（pid == 1 且 kv_out 或 kpe_out 不为空）时不允许早退。
- Hunk 2：将原有的 slot_idx = tl.load(slot_mapping_ptr + tok_idx) 后增加独立判断 if slot_idx < 0: return，用于保留 owner 本地 cache 写操作的负 slot 守卫。
若无法获取 PR 补丁文件，可按评论中提供的 sed/python 脚本方式，对内核文件做正则替换（注意该脚本是针对特定版本的，需核对 hunk 匹配模式）。
重新构建或使用包含补丁的 vLLM 镜像启动服务，保持原有 TP4/DCP4 参数运行 GLM-5.3-NVFP4。

验证方法

以 GLM-5.3-NVFP4 模型在 TP4/DCP4 和 TP4/DCP1 两种配置下分别运行三类提示词：512 个未缓存 token、1536 个未缓存 token、1024 个缓存 + 1536 个未缓存 token；启用 greedy decoding、chunked prefill、prefix caching、FP8 KV cache 和 CUDA graphs，对比输出 token 的 SHA256 哈希。修复后 DCP4 的哈希应与 DCP1 完全一致（Issue 验证中三种场景的哈希分别为 222258a1...、76643708...、4a150ea6...）。也可运行 fused norm/RoPE 内核的完整回归测试（Issue 验证中为 67/67 通过）。