[DeepSeek-V4] Fix nvcc 13 crash building the topk_v2 kernel - #32910
Conversation
topk_small_batch_kernel assigned a cluster-mapped DSMEM address to problem.out, and the problem_transform epilogue loads through that same pointer. cicc segfaults on that merge on CUDA 13.x, so the whole dpsk_v4_topk_v2 JIT module fails to build and a DeepSeek-V4 server on a Hopper or Blackwell host dies at CUDA graph capture with "ninja exited with status 139". Keep the mapped alias in a copy that is only handed to Cluster::forward, which takes TopKProblem by value anyway. The elected rank reads the same bytes back through its own topk_indices, like the Register4 and Streaming branches beside it, so behaviour is unchanged. Adds test_topk_v2_compiles_for_sm90a, which runs nvcc -ptx -arch=sm_90a on the translation unit load_jit() generates. It exits 139 on the pre-fix source and passes on the fixed one, needs a toolkit rather than a cluster-capable GPU, and skips when nvcc is unavailable.
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
|
CUDA 13.2,sglang version above 0.5.14 can not work, #32830, not only DEEPSEEK-V4, glm5.2 has same problm, this fix,can solve it? |
|
Probably yes, if your GLM-5.2 failure is the same nvcc segfault. The kernel this patch fixes Can you confirm your log shows that same failure? The signature is a A quick way to tell without building the patch: run with Fair warning on this PR: I verified the crash and the fix at compile level with nvcc 13.2.86 |
| def test_topk_v2_compiles_for_sm90a(tmp_path) -> None: | ||
| """The kernel must compile for Hopper, whatever the local GPU is. | ||
|
|
||
| cicc crashes on CUDA 13.x if the fused small-batch kernel lets a | ||
| shared::cluster (DSMEM) address reach the pointer its epilogue loads, so the | ||
| whole module fails to build on a Hopper/Blackwell host (issue #32830). Only a | ||
| toolkit is needed to catch that, not a cluster-capable GPU. | ||
| """ | ||
| from tvm_ffi.libinfo import find_dlpack_include_path, find_include_path | ||
|
|
||
| from sglang.kernels.jit.utils.compile import DEFAULT_INCLUDE, KERNEL_PATH | ||
|
|
||
| nvcc = _nvcc_path() | ||
| if nvcc is None: | ||
| pytest.skip("nvcc not found") | ||
|
|
||
| # Same translation unit load_jit() feeds to nvcc for the topk_v2 module. | ||
| source = tmp_path / "topk_v2.cu" | ||
| source.write_text( | ||
| f'#include "{KERNEL_PATH / "csrc" / "deepseek_v4" / "topk_v2.cuh"}"\n' | ||
| "TVM_FFI_DLL_EXPORT_TYPED_FUNC(topk_transform, (TopKKernel::transform));\n" | ||
| "TVM_FFI_DLL_EXPORT_TYPED_FUNC(topk_plan, (TopKKernel::plan));\n" | ||
| ) | ||
| includes = [*DEFAULT_INCLUDE, find_include_path(), find_dlpack_include_path()] | ||
| result = subprocess.run( | ||
| # Device-only (-ptx) keeps this to a few seconds; the flags mirror | ||
| # get_default_target_flags() for a 9.0a target. | ||
| [nvcc, "-ptx", "-arch=sm_90a", "-DSGL_CUDA_ARCH=900"] | ||
| + ["-std=c++20", "-O3", "--expt-relaxed-constexpr"] | ||
| + [f"-I{path}" for path in includes] | ||
| + [str(source), "-o", str(tmp_path / "topk_v2.ptx")], | ||
| capture_output=True, | ||
| text=True, | ||
| ) | ||
| assert ( | ||
| result.returncode == 0 | ||
| ), f"nvcc exited with {result.returncode}\n{result.stderr}" | ||
|
|
DarkSharpness
left a comment
DarkSharpness
left a comment
There was a problem hiding this comment.
LGTM. nvcc crash is really sick 😅
|
I test it，SGLANG_OPT_USE_TOPK_V2=0 can solve this glm5.2 nvcc problem |
Per review, the nvcc-13 cicc segfault is too much of a corner case to keep a dedicated regression test for. test_topk_v2.py goes back to its pre-PR contents; the topk_v2.cuh fix is unchanged.
|
Thanks for the review. Dropped @lucashaha thanks for confirming |
|
/rerun-test test_topk_v2.py |
|
Results for 🚀 |
|
The red CI on this head is infrastructure, not the diff. Could someone kick off a rerun? Only one job failed on its own, That same job also got a 503 from codeload.github.com pulling Everything else is the cascade. The 11 other red jobs ( The two No change on my side, the diff is still just the |
…3.1+
`TopKCluster::forward` selected its phase-3 scatter destination up front:
const auto cur_out = is_primary ? problem.out : smem->tmp_out;
`problem.out` can be a `shared::cluster` (DSMEM) alias of the elected rank's
buffer, while `tmp_out` is `shared::cta`. Merging both into one pointer variable
makes cicc 13.1+ mis-lower the block-local arm for sm_90a and silently drop
every non-primary rank's staged output: `tmp_out` stays zero, and phase 3.5 then
faithfully copies zeros to perfectly correct DSMEM addresses.
The result is a top-k row where only the primary's slots and the `handle_tie`
tail hold valid indices. Slot-diffing one row against a CUDA 12.9 build shows
exactly one contiguous bad run -- [61, 443] -- with slots 0-60 (the primary's own
scatter) and 444-511 (handle_tie) correct. Downstream sparse attention then
dereferences the garbage slots, which is the illegal memory access reported in
sgl-project#33835.
Fix: keep the two destinations in separate code paths so neither pointer ever
carries two address spaces, and note it so nobody merges them back.
Scope: affects any fused small-batch cluster shape (`batch <= 30` and
`seq_len > cluster_floor`), not just rows near the 32K small-batch floor -- the
floor only makes the path reachable at `batch <= 15`. The persistent-pool path
was never affected because it stages output in global memory.
Toolchain matrix (H200, sm_90a, 157-row suite over the report's shapes plus a
boundary sweep and register/streaming/persistent controls):
nvcc before after
12.9 0 bad 0 bad
13.0 0 bad 0 bad
13.1 157 bad 0 bad
13.2 157 bad 0 bad
13.3 157 bad 0 bad
CI builds cu130, which is a clean cell -- that is why this stayed hidden and why
the report came from a CUDA 13.1 deployment.
Also drops the `peer_problem` copy from sgl-project#32910 and states the
block-local pointer at the read-back site instead. That `__builtin_assume` is
load-bearing: removing it reproduces the sgl-project#32830 cicc segfault on 13.1/13.2/13.3.
Validation: 0 bad rows on 12.9/13.0/13.1/13.2/13.3; 1500-iteration randomized
stress and 500 CUDA-graph replays clean; builds for sm_90a and sm_100a on every
toolchain. No performance regression -- worst case +0.1% (noise), and the fused
cluster shapes get 1.3-2.9% faster. sm_100a is compile-verified only; no
Blackwell was available to run on.
…ect#32910) Co-authored-by: guptaishaan <guptaishaan@users.noreply.github.com> Co-authored-by: Xiaoyu Zhang <1182563586@qq.com>
…ect#32910) Co-authored-by: guptaishaan <guptaishaan@users.noreply.github.com> Co-authored-by: Xiaoyu Zhang <1182563586@qq.com>
4 participants
Fixes #32830
topk_small_batch_kernelput a cluster-mapped DSMEM address(
cluster.map_shared_rank(topk_indices, worker_rank)) intoproblem.out, and theproblem_transformepilogue then loads through that same pointer. On CUDA 13.x, ciccsegfaults on that merge, so the whole
dpsk_v4_topk_v2JIT module fails to build and everyDeepSeek-V4 server on a Hopper or Blackwell host with CUDA 13 dies during CUDA graph capture
with
ninja exited with status 139.The mapped alias now lives in a copy that is only passed to
Cluster::forward, which takesTopKProblemby value anyway. The elected rank (blockIdx.y == worker_rank) reads the samebytes back through its own
topk_indices, exactly like the Register4 and Streaming branchesnext to it, so there is no behaviour change.
Verified:
on the exact translation unit
load_jit()generates, targeting sm_90a: same two nvccwarnings, same
Segmentation fault (core dumped), exit 139. Bisected it to the DSMEMpointer plus the later load: removing either one alone also stops the crash, and
-Xcicc -O1compiles the unpatched file, so it is a cicc optimizer crash, not a sourceerror.
test_topk_v2_compiles_for_sm90aintest/registered/kernels/ops/attention/test_topk_v2.pyruns
nvcc -ptx -arch=sm_90aon that translation unit. It fails with exit 139 on thepre-fix source and passes on the fixed source. It needs a toolkit, not a cluster-capable
GPU, and skips when nvcc is not on PATH or under CUDA_HOME.
Not verified: nothing ran on a GPU. The hardware available for this fix was A40 (sm_86),
where
topk_v2.cuhdoes not compile at all (__cluster_dims__ is not supported for this GPU architecture), so the 244 existing correctness cases intest_topk_v2.pycould not be runbefore or after. Please run them on Hopper or Blackwell. I also did not compile the rest of
the JIT tree against CUDA 13.2, so there may be further cicc crashes later in startup; no
other JIT kernel uses
map_shared_rank, so none has this specific trigger.Thanks to @ramon-garcia for the report and the full startup log, which pinned the failure to
the topk_v2 JIT build.
CI States
Latest PR Test (Base): 🚫 Run #30772568799
Latest PR Test (Extra): 🚫 Run #30772573006