Prhub

#32910 [DeepSeek-V4] Fix nvcc 13 crash building the topk_v2 kernel

原始 PR 作者 guptaishaan 合并时间 2026-08-03 10:59 文件变更 1 提交数 3 评论 9 代码增减 +7 / -2

bugfix jit-kernel deepseek blackwell

执行摘要

修复 nvcc 13 下 topk_v2 JIT 内核编译崩溃

issue #32830 报告：用户以 CUDA 13.2 在 4 卡 H100 上启动 DeepSeek-V4-Flash 时，nvcc 编译内核阶段直接段错误退出。PR body 进一步定位：topk_small_batch_kernel 把 cluster.map_shared_rank(topk_indices, worker_rank) 返回的 DSMEM 地址写入 problem.out，而 problem_transform epilogue 经同一指针加载数据，cicc 在 CUDA 13.x 上对该合并模式段错误，导致整个 dpsk_v4_topk_v2 JIT 模块构建失败，"every DeepSeek-V4 server on a Hopper or Blackwell host with CUDA 13 dies during CUDA graph capture with ninja exited with status 139"。issue 评论区进一步确认 GLM-5.2 同样受影响，即该内核为 DSV4 与 GLM DSA 共享，修复具有跨模型收益。

值得阅读。改动本身仅 7 行，但 PR body 完整记录了一次教科书式的编译器崩溃诊断：在 sglang 外独立复现 load_jit() 生成的翻译单元、二分定位触发模式、用 -Xcicc -O1 证明是优化器缺陷、并跨 nvcc 12.6 / 13.2 验证兼容性。对需要维护 JIT kernel 与多 CUDA 工具链的工程师有直接参考价值；"以按值传参副本隔离仅供写入的别名指针"的做法也是规避编译器 bug 而不改变语义的干净范例。建议关注合并后 Hopper / Blackwell 上的正确性回归结果。

讨论亮点
  1. 测试取舍：DarkSharpness 对新增的 test_topk_v2_compiles_for_sm90a 提出质疑："I don't think we should include this test for this extremely corner case. What do you think @BBuf @zcnrex"。作者随后把 test_topk_v2.py 恢复原状，DarkSharpness 在最终 review 中 APPROVED 并评价 "LGTM. nvcc crash is really sick 😅"。
  2. GLM-5.2 波及确认：issue 评论区 lucashaha 报告 GLM-5.2 同样在 CUDA 13.2 下无法启动，并实测 SGLANG_OPT_USE_TOPK_V2=0 可绕过；guptaishaan 确认 topk_v2.cuh 是 DSV3.2 / DSV4 与 GLM DSA 共享的内核（python/sglang/srt/layers/attention/dsa/dsa_topk_backend.py），修复落地后应无需环境变量，并建议用户若仍失败再 ping 回报。
  3. CI 红色归因：guptaishaan 举证 11 个红 job 均为 wait-for-base-b fast-fail 级联，根因是 GitHub 侧下载 sglang-kernel==0.4.5+cu130 wheel 与 actions/checkout 的 503 故障，与 diff 无关；随后 /rerun-test test_topk_v2.py 在 1-gpu-h100 上绿色。

实现拆解

  1. 根因定位：作者在 sglang 外以报告者的工具链（nvcc 13.2.86 + gcc 13.4）对 load_jit() 生成的精确翻译单元做 sm_90a 编译复现，得到一致的两个警告与 exit 139；随后二分确认崩溃来自 DSMEM 指针与后续 load 的组合，且 -Xcicc -O1 可以编译未修补文件，锁定为 cicc 优化器缺陷而非源码错误。
  2. 核心修复：python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh 中 topk_small_batch_kernel 的 cluster 分支（blockIdx.y != worker_rank 侧），把原来的"直接对 problem.out 赋映射地址再调 Cluster::forward"改为：先浅拷贝 peer_problem = problem，只对 peer_problem.out 赋映射地址，再按值传入 Cluster::forward。因 Cluster::forward 本就按值接收 TopKProblem，被选中的 rank 仍通过自己的 topk_indices 读回同一份字节，与相邻的 Register4 / Streaming 分支语义一致，无行为变化。
  3. 兼容性验证：补丁后同一翻译单元在 nvcc 13.2.86 与 12.6 下均编译通过；无 GPU 运行验证（作者仅有 A40 / sm_86，该内核在该架构下根本不支持 cluster_dims），244 个既有正确性用例未能在本地执行，PR body 明确请求在 Hopper / Blackwell 上补跑。
  4. 测试配套演进：初始提交新增 test_topk_v2_compiles_for_sm90a（用 nvcc -ptx -arch=sm_90a 编译该翻译单元，无需集群 GPU，缺 nvcc 时跳过），经 DarkSharpness review 认为 corner case 不值得保留专门回归测试后删除，test_topk_v2.py 恢复 PR 前内容；/rerun-test test_topk_v2.py 在 1-gpu-h100 上通过。
文件 模块 状态 重要度
python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh JIT 内核 modified 3.93

关键符号

topk_small_batch_kernel

关键源码片段

python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh core-logic

唯一变更文件：topk_small_batch_kernel 的 cluster 分支将 DSMEM 映射指针移入 TopKProblem 副本，规避 CUDA 13.x cicc 优化器段错误，解除 DSV4 / GLM DSA 在 Hopper / Blackwell 上的启动崩溃。

// 整理自 python/sglang/kernels/jit/csrc/deepseek_v4/topk_v2.cuh 中的
// topk_small_batch_kernel 函数，展示其 cluster 协作分支（修复后版本）。
// 前置逻辑已省略，包括 smem 初始化、problem 填充、worker_rank 计算等。
if (blockIdx.y == worker_rank) {
    // 本 rank 直接走 streaming 分支，读取自己的 topk_indices
    Streaming::forward<kPDL>(problem, &smem);
} else {
    auto cluster = cooperative_groups::this_cluster();
​
    // 修复说明：map_shared_rank 返回 shared::cluster（DSMEM）的映射地址。
    // 旧代码把该地址直接写入 problem.out，而 problem_transform epilogue
    // 会经由同一指针加载数据；cicc 在 CUDA 13.x 上编译这一模式时发生
    // 段错误（exit 139），导致 dpsk_v4_topk_v2 JIT 模块构建失败，DSV4
    // 与 GLM DSA 服务在 Hopper / Blackwell + CUDA 13 环境启动即崩溃
    // （issue #32830）。修复方式是把映射别名放进 TopKProblem 的浅拷贝
    // peer_problem，仅将其传给按值接收参数的 Cluster::forward；被选中的
    // rank（blockIdx.y == worker_rank）仍通过自己的 topk_indices 读回
    // 同一份字节，运行语义完全不变。
    auto peer_problem = problem;
    peer_problem.out = cluster.map_shared_rank(topk_indices, worker_rank);
    Cluster::forward<kPDL>(peer_problem, &smem); // 写入 peer 的输出共享内存
    cluster.sync();
}

评论区精华

是否保留 sm_90a 编译回归测试 测试

DarkSharpness 对新增的 test_topk_v2_compiles_for_sm90a 提出质疑，认为这是极端 corner case，不值得保留专门回归测试，并询问 @BBuf @zcnrex 意见。

结论：作者同意并删除该测试，test_topk_v2.py 恢复 PR 前内容；DarkSharpness 随后 APPROVED 并评价 LGTM。 · 已解决

GLM-5.2 是否同样受影响及绕过方式 question

lucashaha 报告 GLM-5.2 在 CUDA 13.2 下同样无法启动，询问本修复能否解决；guptaishaan 解释 topk_v2.cuh 是 DSV3.2 / DSV4 与 GLM DSA 共享内核，建议用 SGLANG_OPT_USE_TOPK_V2=0 验证；lucashaha 实测确认可绕过。

结论：确认同一内核、同一 cicc 崩溃，修复应覆盖 GLM DSA；作者未亲自在 GLM-5.2 上验证，但用户侧已确认绕过路径。 · 已解决

红色 CI 是否由 diff 引入 other

guptaishaan 举证 11 个 red job 均为 wait-for-base-b fast-fail 级联，根因是 GitHub 侧下载 sglang-kernel wheel 与 actions/checkout 的 503 故障，发生在任何 sglang 代码运行之前。

结论：确认为基础设施故障，非 diff 引入；/rerun-test test_topk_v2.py 在 1-gpu-h100 上通过。 · 已解决

风险与影响

  1. 缺少目标硬件回归：244 个正确性用例未在 Hopper / Blackwell 上运行，作者论证逻辑等价（被选中的 rank 读回同一份字节），但 cluster 协作路径与 DSMEM 读写仍需在目标架构确认。
  2. 回归测试已移除：该 cicc 崩溃点无自动化防护，未来若重构 topk_small_batch_kernel 的拷贝逻辑可能重新引入且 CI 无法发现。
  3. JIT 树未全量验证：作者未对 CUDA 13.2 编译其余 JIT 内核，启动路径上仍可能存在其他 cicc 崩溃；不过没有任何其他内核使用 map_shared_rank，此特定触发器唯一。
  4. 隔离依赖既有约定：修复正确性依赖 Cluster::forward 按值接收 TopKProblem；若未来改为引用传递，副本隔离将失效。

对 CUDA 13.x + H100 / B200 等 Hopper / Blackwell 用户是启动即崩的阻断性问题的解除，直接受益模型为 DeepSeek-V4 与 GLM DSA 系列；对 CUDA 12.x 用户及已用 SGLANG_OPT_USE_TOPK_V2=0 绕过的用户无行为变化。改动仅限单文件 7 行，JIT 编译产物与运行时行为不变，风险面小。团队层面沉淀了一套可复用的编译器崩溃诊断范式（独立翻译单元复现 + 二分 + -Xcicc -O1 缓解实验），对后续 CUDA 13 全面升级有参考价值。

CUDA 13 编译器兼容隐患 缺少 Hopper/Blackwell 硬件回归 回归测试被移除 JIT 树其余内核未验证

关联 Issue

#32830 [Bug] Nvidia compiler crashes with segmentation fault when trying to serve DeepkSeek v4

完整报告

Markdown PDF

参与讨论