[Bug] XPU qnorm/rope kernel: overlapping stores in one program corrupt the KV cache nondeterministically

快速结论：这个报错通常出现在使用 vLLM 的 XPU 版 DeepSeek V4 融合 qnorm/rope/KV 写入算子（xpu_qnorm_rope_kv_fp8_insert）时，KV cache 的 RoPE 区域被同一 program 内的批量写入和旋转写入相互覆盖，导致缓存结果不确定。优先排查 KV 分支的批量 store 是否已加上 NOPE_DIM 掩码。

适用环境：Issue 已确认存在于 vLLM v0.27.1（tag 6e448d0ea9bf）及当时 main；原始复现环境为 MTT S5000（MUSA）；修复验证环境为 Intel Arc Pro B60，镜像 vllm/vllm-openai-xpu:v0.32.0，torch 2.14.0+xpu，Triton 3.8.0。

最快修复方案：在 KV 分支的批量 store 上加入 mask=offs < NOPE_DIM（与 Q 分支已有的 nope_mask 写法一致）。该修复由 PR #46685 合入 main，Issue 报告者验证后确认可关闭为已修复。

注意事项：XPU 后端调度可能不会暴露该问题，但“两次写入同一地址且无顺序保证”在任何后端都不成立，不能因 XPU 上未复现就忽略；仅升级到带 #46685 的版本即可，无需其他配置改动。fp8 量化带来的少量舍入差异属于正常现象，与覆盖竞态无关。

问题场景

用户在 vLLM 中运行 DeepSeek V4 的 XPU 融合算子 vllm/models/deepseek_v4/xpu/xpu_qnorm_rope_kv_fp8_insert.py，执行 KV 分支写入时触发。以 4096 tokens × 64 heads 的规模，将结果与 FlagGems 融合算子（该算子在相同输入上与纯 torch 参考实现逐字节一致）对比，发现部分 token 的 KV cache RoPE 区域内容不正确，且两次相同输入运行的结果不同。原始发现在 MTT S5000（MUSA）上，因为 MUSA 的调度方式不同而暴露。

报错原文

[Bug] XPU qnorm/rope kernel: overlapping stores in one program corrupt the KV cache nondeterministically

原因分析

KV 分支存在同一 program 内对相同地址的两次写入，且没有顺序保证：

# lines 93-94
kv_full = tl.load(kv_base + offs)
tl.store(kv_out_base + offs, kv_full)        # 写入未旋转的 rope 区域
...
# lines 106-107
tl.store(kv_out_base + even_offs, new_even)  # stride-2 地址覆盖
tl.store(kv_out_base + odd_offs, new_odd)

两次 store 的 lane 到地址映射不同，中间没有同步或屏障，无法保证旋转后的值最终落地。证据显示，在 4096 tokens 中有 30 个 token 的缓存 RoPE 区域不一致，其中 29 个恰好等于未旋转的原始 kv 值，说明批量写入最后生效；另 1 个为部分覆盖，符合 lane 粒度的竞态特征。第二次相同输入运行得到 29 个而非 30 个不一致 token，说明结果不确定。Q 分支因为已经使用 nope_mask 屏蔽了 NOPE 之外的地址，所以从未出现该问题。

环境排查

确认 vLLM 版本：v0.27.1（tag 6e448d0ea9bf）及当时 main 均存在该问题。
确认是否已包含 PR #46685（合入 main，commit 6463d04601），或 v0.32.0 镜像等更新的构建。
确认 GPU/加速卡类型（XPU、MUSA 等）及 Triton 版本，例如验证环境为 Triton 3.8.0、torch 2.14.0+xpu。
确认触发规模：小规模（如 16 tokens × 1 head）可能不暴露，256 × 8 起开始复现，2048 以上基本每次运行都会出现差异。

解决步骤

确认当前 vLLM 源码中 KV 分支的批量 store 是否为未加掩码的 tl.store(kv_out_base + offs, kv_full)。
升级到已包含 PR #46685 的版本（main 或后续发布版本）；若需本地修补，将该行改为 tl.store(kv_out_base + offs, kv_full, mask=offs < NOPE_DIM)，与 Q 分支第 70 行的 nope_mask 模式保持一致。
该掩码同时省去每个 KV token 64 次冗余写入，属于等价且更省的做法。
若自行打补丁，建议参照 PR 增加针对 KV RoPE 区域的回归测试（问题讨论中已随修复添加）。

验证方法

在相同输入上多次运行该算子（建议 20 次），将第 2 至第 20 次输出与第 1 次逐字节比较：修复后 KV cache 与 q 均应在所有形状上零差异。验证者给出的结果为 256×8、2048×8、8192×8、8192×64 均 0/19，而未加掩码的对照版本在 256×8 时 9/19、2048 以上 19/19 不一致。另外可将 cache 与 torch 参考路径（RoPE 后走同一 quantize_and_insert_k_cache）对比，残余差异应仅为少量 fp8 舍入（最多 2 LSB，且不落在 scale 字节），并在每次运行中保持一致。