Conversation
FullCudaGraphBackend.capture_one() synchronizes the device before each of the two warmup iterations but not between the last warmup and torch.cuda.CUDAGraph() creation. In compact ragged-verify mode, each new non-uniform verify_lens shape triggers fresh JIT compilation of compressor/norm_rope kernels during warmup. If that compilation is still in flight when capture begins, the JIT issues CUDA driver calls (cuModuleLoadData, etc.) inside the capture context, which is illegal and surfaces as CUDA_ERROR_ILLEGAL_ADDRESS or cudaErrorStreamCaptureUnsupported. This predominantly affects compact mode where bs values with uniform verify_lens reuse previously compiled kernels but the first non-uniform shape (e.g. bs=60 => [3]*104+[2]*24) compiles new kernels right as capture starts. static/cap-accept modes are unaffected because they use uniform verify_lens for all shapes, so kernels are compiled once on the first shape. Verified: without the sync, compact capture crashes deterministically at bs=60 (3/3 runs, H200x4 TP4 DSV4-Flash-0731); with it, all 35+35 shapes capture and the server reaches ready. Fixes the root cause reported in sgl-project#33356.
…test Regression guard for sgl-project#33356 / PR sgl-project#33795. Pins the ordering invariant in FullCudaGraphBackend.capture_one(): a device synchronize + TP-group barrier must run between the last warmup forward and torch.cuda.CUDAGraph() construction, so async TVM/DeepGEMM JIT cannot issue driver calls inside the stream-capture region. Deterministic mock-only call-order check - no GPU.
|
Two updates:
|
|
We have a third-environment data point on the "Related #32467 — So at least one environment in the #33356 failure class appears not to be Happy to run this PR's patch on our deterministic repro with and without |
|
Ran the arms we offered — full matrix and evidence in #33356. Two results relevant to this PR: Your patch composes cleanly with #32467 (arm D). With both applied, our SM120 box captures the full ladder at the deterministic-fault shape and at the worst row of our bisection matrix (ctx 1M, mem 0.90, cap 64, expandable segments on — 27 shapes in 72.9 s), with normal capture times. No interaction, no regression. That supports your own "complementary, land both" note. Your patch alone (arm C) does not cover our failure. The traceback lands at That's the basis for the earlier suggestion about |
|
@hassellof — thanks, the 4-arm matrix is exactly the separation this issue needed. The arm C traceback at full_cuda_graph_backend.py:92 (forward_fn() in the warmup loop, before the added sync at 96-105 and before CUDAGraph() at 107) cleanly separates the two producers:
I've updated the PR wording to Related #33356 so the issue stays open for the P2/SM120 case until #32467 lands (it has been open since 07-27 and on your matrix it closes the deterministic instance — agreed it should land alongside this one; arm D shows they compose with zero interaction). We'll add H200-side confirmation of arms C and D to the thread (our bs=60 is a deterministic P1 repro; ~30 min per arm) so the landing case has both environments covered. Happy to run any specific capture-ladder or allocator row you want on the H200 box. |
|
Could a maintainer approve the workflow runs for this fork PR? All 118 jobs are currently skipped pending CI authorization, including base-a-test-cpu which executes the new regression test (test_full_cuda_graph_backend.py) added at head 994910f. The change is 2 files (+125): the capture_one() sync/barrier fix and its deterministic mock-only unit test. Happy to rebase/address anything else needed for review. |
|
Great — the P1/P2 split reads right to me, and thanks for switching to Most useful H200 rows, if you're offering. The two that carry the most information are the ones where our boundary moved with allocator config rather than with ladder depth alone:
Rows 1→2 are the depth boundary; 1→3 is the same-or-shallower ladder faulting once expandable segments are on; row 4 is the odd one — a single 240 tier is clean while a ladder containing 240 faults, which we could never fully separate from a context co-factor. If P1 is your only producer on H200, I'd expect all four to behave identically there with your sync applied and #32467 absent (arm C), since none of them should be reachable by a device-side Offer in return: we have a 4× RTX PRO 6000 (sm_120) box available and our compact-mode repro is deterministic, ~30 min per arm. Happy to run your bs=60 P1 repro on it, or any instrumented build — if P1 turns out to be reachable on SM120 too, that's worth knowing before either patch lands. On CI: we're stuck at the same gate on #33407, #33813 and #33816 (all jobs skipped pending the |
Independent B300/TP8 validation: the patch is insufficient on the original #33356 environmentI tested this PR on the original B300-class reproducer from #33356, using the same container/task, machine, model, and capture configuration throughout. EnvironmentThe PR patch was applied as the exact single hunk from Controlled results
Aggregate for the patched configuration: Given the pre-existing non-determinism, the single READY launch is not sufficient evidence that the patch fixes this environment. The failing shape, rank, and surfacing path continued to drift after the patch: Therefore, on B300 x8 / TP8 / DeepSeek-V4-Pro-DSpark / v0.5.16-cu130, the final device sync + TP barrier is insufficient by itself. I do not consider #33356 validated or resolved by this patch, and the issue should remain open. DeepGEMM precompile controlI also ran All three subsequent launches still printed The remaining failure is still consistent with some asynchronous activity or lifetime violation near the capture boundary, but these data do not identify the original producer. Candidate areas remain first-use JIT/module loading not covered by the precompile sweep, MoE-routing workspace lifetime, CUDA-Graph pool address reuse, and alternate-stream ownership/lifetime. Full local evidence includes seven logs, the capture-results table, before/after source hashes, runtime import/order verification, and the exact applied diff. |
Follow-up: completed #32467 × #33795 four-arm matrix on B300/TP8This supplements my earlier B300 result. I completed the complementary-kernel-fix matrix using the same original #33356 environment. Fixed environmentResults
All six runs with #32467 enabled were free of illegal-memory, SIGSEGV, capture-failure, and scheduler-initialization signatures. The source state was checked before each arm: The TVM-FFI JIT cache was cleared after changing Updated interpretationFor the original B300/TP8 startup-capture failure, #33795 alone is not sufficient (
This does not invalidate the reported H200 result or the general warmup-to-capture ordering invariant; it narrows the B300 root cause and shows that the drifting Triton/DeepGEMM/TRT-LLM surfacing kernels were downstream symptoms of the plan-kernel race. I also tested Scope remains startup capture only; sustained runtime replay from #31023 is separate. |
Follow-up: four-arm matrix with #32467 on B300 / TP8Following my earlier B300 result where #33795 alone reached READY in only 1/6 fresh runs, I completed a controlled EnvironmentMatrix
Arm B explicitly restored the unpatched InterpretationFor this B300/TP8 reproducer:
This indicates that the primary failure here is the plan-kernel race fixed by #32467, rather than the final warmup-to-capture synchronization gap. This does not invalidate the H200 evidence or the ordering guard added by this PR; it means #33795 is not the necessary recovery condition for this specific B300 case once #32467 is present. I also attempted |
|
Thanks for the concrete rows — very helpful. Here's what we have so far, and where we can fill in. H200×4 / TP4 / DeepSeek-V4-Flash-0731 — arm D (both #33795 + #32467) We ran a 4-stage validation on a dedicated profiling pod (image
BCG is architecturally disabled for DSV4 ( We also ran a #31023 stress test (128 concurrent, 8192 input / 1000 output tokens, 1280 requests, compact + graph + overlap ON): all completed in 572s, zero IMA. Bug 2 did not reproduce on TP4 H200 (original repro was TP8 B300). What we haven't run yet — arm C on H200 Your 4-row request is exactly the right experiment. We have not yet run arm C (#33795 only, #32467 absent) on H200. Our image has both patches baked in, so we'd need to rebuild with only #33795 to test this. If P1 is the only producer on H200, all 4 rows should pass with #33795 alone; if any faults, P2 is reachable on Hopper too. Based on @Phoenix3334's B300 matrix just posted (arm C = 1/6 READY on B300×8 TP8), the picture is now:
The open question is whether H200×4 is P1-only (in which case #33795 is necessary there) or also has a P2 component (in which case #32467 alone might suffice there too, matching B300). We'll run arm C with your 4 rows and report back — should take ~2h to rebuild the image and run the matrix. On the landing question: given that #33795 alone is insufficient on 2 of 3 environments, and #32467 alone is sufficient on those same 2, the case for landing #32467 first (or at least simultaneously) is strong. #33795 remains a correct ordering guard regardless — the post-warmup |
|
Follow-up from the B300/TP8 four-arm matrix: I also completed runtime replay validation with #32467 ON and this PR OFF. The startup matrix remains: Runtime validation of Arm B (#32467 only):
A second fresh Gate 2' run was manually stopped at 293/512 to release the machine; it had no GPU error signature up to that point and is not counted as pass/fail. So on this B300x8/TP8 reproducer, #32467 alone is sufficient for both startup capture and sustained replay in the tested paths. This PR still composes cleanly with #32467 (Arm D 3/3), but it was not required for the successful B300 startup or replay runs above. This does not contradict the H200 result reported for #33795; it narrows the B300 failure class tracked in #33356 toward the plan-kernel race fixed by #32467. |
|
One more B300 data point relevant to the interaction with #32467: I now reproduced the #32467 failure mechanism directly in a tiny single-GPU With the unpatched The bad maximum ID is exactly That makes the B300 distinction between the two fixes clearer:
This is consistent with our earlier four-arm result (#33795-only 1/6 READY vs #32467-only 3/3 READY) and sustained-replay result (#32467-only 320/320 and 512/512 clean). I still view the two changes as orthogonal/compatible; this update only narrows the B300 root cause to the plan race fixed by #32467. |
|
One more data point that sharpens the scope distinction here: on the B300 failure class from #33356, I can now reproduce the corruption without CUDA Graph capture at all, by directly exercising the real GPU JIT With #32467 removed: With #32467 applied and the TVM-FFI cache cleared/recompiled: For both failing ragged cases, So for the B300 branch of #33356, the corruption is demonstrably a kernel-plan correctness race fixed by #32467, not a capture-context-only failure. That does not invalidate the H200 behavior motivating this PR; it does mean the statement in this PR body that the #33356 failure is "a capture-context violation, not a kernel bug" is too broad across environments. The four-arm B300 matrix remains consistent with this:
and #32467-only also passed sustained replay gates afterward. I would keep #33795 scoped as the H200 capture-ordering fix / complementary fix rather than the root cause for the B300 case. |
|
One more B300 data point that sharpens the interaction with #32467: I now have a deterministic single-GPU kernel-level reproducer for the Directly calling the real For both failing ragged cases, That directly demonstrates the #32467 mechanism ( So for the B300/TP8 case, this is stronger evidence that the dominant failure is the kernel correctness race rather than the post-warmup capture-ordering gap fixed here. This does not invalidate the H200 ordering issue addressed by #33795; it just separates the two mechanisms more cleanly. |
Follow-up from the B300 investigation: #32467 now has a direct kernel-level reproducerI reduced the B300 failure to the real single-GPU JIT Unpatched #32467: With #32467 applied: For both failing ragged cases the bad maximum was exactly which is the exact upper bound of the MTP-uniform mapping appearing on a non-uniform batch. This directly demonstrates the missing- So the B300/TP8 evidence now separates the two fixes more cleanly:
This does not invalidate the H200 ordering evidence behind this PR; it only shows that the original B300 failure had an independently demonstrated device-kernel root cause. Full reduced-test details are posted on #32467 and the consolidated status is in #33356. |
Follow-up from B300: deterministic kernel reproducer points to #32467, not capture orderingOne more data point after the earlier B300 four-arm matrix. I reduced the failure to a single-GPU harness around Results:
The bad ragged plans reach This makes the B300 four-arm result easier to interpret:
Runtime under Arm B also passed:
So for this B300/TP8 reproducer, the producer-side bug is now directly demonstrated in the plan kernel and fixed by #32467. The final device-sync + TP-barrier change in this PR remains compatible, but it is not required to make the B300 case stable. This does not invalidate the H200 result that motivated #33795; it indicates that at least two distinct races/signatures can exist across platforms, and the B300 case was dominated by the plan-kernel race fixed in #32467. |
…project#32035 sgl-project#33656 sgl-project#32183 sgl-project#33145) Applied PRs (latest from GitHub): sgl-project#33288 Indexer logits OOM fix sgl-project#30393 HiCache packed/sidecar draft caches sgl-project#31170 DPA prefix_affinity load balancing sgl-project#33795 DSpark compact ragged-verify CUDA graph JIT race sgl-project#32467 C128 plan-kernel warp barrier sgl-project#33865 DSpark x prefill CP unblock sgl-project#30371 SWA state pool sizing (storage page) sgl-project#33358 FlashMLA norm-rope K-tokens-per-block ILP sgl-project#33872 num_draft_tokens clamp + extend_len==0 skip (supersede sgl-project#32183) sgl-project#34002 Sidecar backup vacuously-successful fix (replaces sgl-project#33656, with tests) sgl-project#33862 Reclaim redundant host mirrors after storage backup sgl-project#31315 Avoid repeated Mooncake gets after stale hits sgl-project#32327 Q8KV8 sparse MLA prefill backend (flashmla_sparse_q8) sgl-project#31668 Fix sidecar pool life-time (use-after-free on prefetch abort) sgl-project#31195 TP0 verify-token-budget broadcast (adapted to get_schedule() API) Dropped (per user request or superseded): sgl-project#32771 IndexCache C4 top-k reuse — has bug sgl-project#32035 DSpark C128 online compressor — has bug sgl-project#33656 Superseded by sgl-project#34002 (same fix + unit tests) sgl-project#32183 Superseded by sgl-project#33872 (included in supersede PR) sgl-project#33145 Base f01f706 already has superior reasoning-effort profile system Conflicts resolved: sgl-project#31195: adapted to base get_schedule().disable_overlap_schedule API sgl-project#32327: path remapped jit_kernel/ -> kernels/jit/ and kernels/ops/attention/ sgl-project#31668: applied cleanly on top of sgl-project#30393+sgl-project#34002+sgl-project#33862 modifications
B30Z/TP8 v0.5.17 follow-up after the H200
|
| #32467 | this PR (#33795) | max decode graph bs | Result |
|---|---|---|---|
| OFF | OFF | 128 | 1/1 capture failure at bs=104, TP7 (CUBLAS_STATUS_EXECUTION_FAILED) |
| ON | ON | 128 | 3/3 READY, 35/35 shapes each |
| ON | OFF | 128 | 1/1 READY, 35/35 shapes |
| ON | ON | 256 | 1/1 READY, 51/51 shapes |
All five patched runs captured bs=10 successfully. The H200-reported v0.5.17 bs=10 failure therefore did not reproduce on this B30Z/TP8 stack.
For this platform the new result is consistent with the earlier v0.5.16 matrix: #33795 composes cleanly with #32467, but it was not required in the one isolated v0.5.17 #32467-only arm. I am deliberately not generalizing that to H200 or calling this PR unnecessary: TP topology, GPU architecture, CUDA/driver/library stack, model/build provenance, and the capture-shape set differ from the H200 report.
So the current evidence supports a narrower statement:
- the warmup-to-capture synchronization invariant in this PR may address a real environment-dependent ordering problem;
- it is not the primary fix for the original B30Z [Bug] DSpark large decode CUDA-Graph capture can hit non-deterministic illegal memory on TP8 (v0.5.16) #33356 failure class, where the producer-side
c_plan.cuhrace from [BugFix] Fix race in c128 prefill plan kernel on ragged extend #32467 has direct OOB-plan evidence; - the additional H200 v0.5.17 failure remains unresolved and is not currently reproducible on our B30Z/TP8 setup.
|
Related backend-level follow-up: the same structural warmup-to-capture completion gap is present in The Breakable ordering before the fix is analogous: #34286 adds the same boundary: and adds a CPU-only deterministic ordering regression requiring: I first ran that regression against the unpatched Breakable backend; it went RED specifically because the post-warmup completion boundary was absent. After the production change it passed, and the relevant runner-backend unit suite was 5/5. I am treating this as evidence for a shared capture-initialization invariant, not as a claim that Breakable CUDA Graph reproduces the H200 DSpark crash from this PR. The B300 producer bug remains separately localized to #32467. The cross-backend contract is tracked in #32432. |
|
2026-08-11 clarification on the relationship to #32467: the final #32467 implementation has changed from the original barrier experiment to the reviewer-selected no-init formulation ( Same-session A/B/C producer regression: The B300 startup isolation also keeps the two mechanisms separate: So for the tested B300 path, #32467 is the primary producer-side fix; #33795 remains a separate FullCudaGraphBackend capture-initialization ordering fix with its own H200 evidence. The complete current root-cause map and today's TP8/DP8 decode-only runtime data are in #34297. |
Problem
Both CUDA-graph backends run two warmup iterations before capture, with a
`synchronize()` + `barrier()` at the **top** of each iteration:
```python
for warmup_step in range(2):
self._device_module.synchronize()
self._tp_group.barrier()
with self._precarve.measure():
output = forward_fn()
...
if post_warmup_hook is not None:
post_warmup_hook()
... nothing drains the last warmup's own async work ...
graph = torch.cuda.CUDAGraph()
```
Those syncs order work issued *before* each warmup. Nothing drains the **final**
warmup's own async work.
Impact
Lazy / JIT kernel compilation kicked off by a first-seen shape can still be in flight
when capture starts. A JIT module that finishes loading mid-capture issues illegal
driver calls (`cuModuleLoadData` and friends) on the capturing stream, surfacing as
`CUDA_ERROR_ILLEGAL_ADDRESS` or `cudaErrorStreamCaptureUnsupported`.
This is intermittent by nature — it depends on whether the last warmup happened to
trigger a compile and whether that compile lands before or after capture begins.
Fix
Drain the device and re-align ranks once more, immediately before the capture.
Applied to **both** backends. sgl-project#33795 fixed
`FullCudaGraphBackend`; `BreakableCudaGraphBackend` has the identical pattern and was
left unfixed.
Verification
Reasoning-only. The ordering gap is visible in the source and the fix is the same
drain sgl-project#33795 already established as correct for the sibling backend. I have not
constructed a reproducer — the failure is a race and reproducing it reliably would
mean forcing a JIT compile on the last warmup shape.
Cost is two extra synchronization points per captured graph, at capture time only.
No effect on the replay path.
---
*Provenance: originally authored by an automated kernel-optimization agent
(hyperloom session `20260806T082051Z`), rebased onto current `main` and reviewed by
hand before submission.*
3 participants
Motivation
In DSpark compact ragged-verify mode, large decode CUDA-Graph capture can fail during server startup with a non-deterministic illegal-memory access or host SIGSEGV (#33356). Across runs the failure surfaces through different kernels (DeepGEMM, c128_v2, Triton routing, fused_norm_rope_v2) and different TP ranks, which makes it look like a wild kernel bug — but it is a capture-context violation, not a kernel bug.
FullCudaGraphBackend.capture_one()synchronizes the device before each of the two warmup iterations, but there is no synchronization between the last warmup andtorch.cuda.CUDAGraph()creation. Warmup forward passes trigger asynchronous JIT compilation of compressor/plan kernels. In compact mode each new non-uniformverify_lensshape compiles fresh kernels during warmup; if that compilation is still in flight when capture begins, the JIT issues CUDA driver calls (cuModuleLoadData, etc.) inside the capture context, which is illegal — the first kernel to run after that is the one that reports the IMA.This also explains the deterministic shape signature seen on H200×4:
verify_lens(e.g. bs=128..64) reuse the kernels compiled by the previous shape → no new JIT → clean;[3]*104 + [2]*24) makes the planner take the general-prefill path → fresh kernels get compiled right as capture begins → race window opens;static/cap-acceptare unaffected because they use uniformverify_lenseverywhere, so kernels are compiled once on the first shape.Modifications
Add a device
synchronize()+ TP-groupbarrier()after the warmup loop inFullCudaGraphBackend.capture_one(), before entering the capture context:Files changed:
python/sglang/srt/model_executor/runner_backend/full_cuda_graph_backend.py(+11)Accuracy Tests
This change does not alter any kernel or model-forward logic — it only inserts a device sync + TP barrier between warmup and capture. Outputs are bit-identical. No accuracy test is affected.
Speed Tests and Profiling
Setup: H200 × 4 (TP4), DeepSeek-V4-Flash-0731, DSpark (gamma=5,
--speculative-num-draft-tokens 6),SGLANG_RAGGED_VERIFY_MODE=compact,--cuda-graph-max-bs 128,--max-running-requests 128,--moe-runner-backend flashinfer_mxfp4.Startup overhead: one extra
synchronize()+barrier()per captured shape (~70 shapes) — negligible versus capture time. No steady-state inference cost (the sync is only in the one-time capture path).Checklist
Related
plan_compress_prefill_kernel0(__syncthreads()missing between warp-scratch init and min/max reduction) that can mis-select the MTP fast path for ragged extend and generate out-of-boundsragged_id. Both fixes live in the same compact capture path but fix different races; landing both is recommended.CI States
Latest PR Test (Base): ❌ Run #31067800551
Latest PR Test (Extra): ❌ Run #31067800480