Conversation
Signed-off-by: EanWang211123 <wangyiheng@sangfor.com.cn>
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
1 similar comment
|
Caution The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased. |
|
@yhyang201 I see that the |
|
Independent hardware validation for this fix, in case it helps the review @yhyang201 @lmyybh. On 4× RTX PRO 6000 Blackwell (sm_120), TP=4 + DP attention + EP, DeepSeek-V4-Flash-0731, sglang v0.5.16 with this PR backported (paths rewritten for the pre-refactor tree): it fixes a deterministic decode-graph capture IMA under Before the backport we had mapped an 8-row shape/allocator matrix where the fault threshold moved with the CUDA-graph batch-size ladder and with Two notes that may be useful:
We are also running the 4-arm matrix proposed in #33356 (this PR × #33795's capture-ordering fix) on the same hardware and will report the result there. CI has not run on this PR — if a maintainer can add the |
Independent B300/TP8 validation: PASSI tested the EnvironmentThe v0.5.16 source path is __syncthreads();I cleared the TVM-FFI JIT cache before the first fixed run so the modified header could not reuse the old compiled artifact. Four-arm result
Per-run startup times: Across all six runs with #32467 enabled, log scans found no: InterpretationOn this B300/TP8 configuration, the missing block synchronization fixed by #32467 is the main cause of the startup-time compact ragged CUDA-Graph capture IMA tracked in #33356. The post-warmup sync/barrier from #33795 is compatible, but was neither sufficient by itself nor required once this kernel race was fixed. This result is scoped to startup capture ( Tested-by: @Phoenix3334 |
Independent B300 / TP8 validationI validated this fix on the original B300 reproducer from #33356 using a controlled EnvironmentThe v0.5.16 source tree uses the equivalent file at: I applied the same semantic change from this PR: a single Four-arm result
Arm B elapsed times were 331/297/291 s; Arm D was 298/296/296 s. All six fixed-kernel runs completed capture and reached HTTP ready. Scans for Controls
ConclusionFor this B300×8 / TP8 reproducer, #32467 is the primary fix for the startup capture IMA. #33795 alone was insufficient, while #32467 alone was sufficient in 3/3 fresh runs; combining both fixes was also stable in 3/3 runs. This matches the proposed mechanism: the missing block barrier can misclassify non-uniform ragged input as uniform and generate out-of-bounds Tested-by: @Phoenix3334 |
|
Third-hardware datapoint: TP4 H200×4 validation (arm D, both patches) Adding to @hassellof's SM120 and @Phoenix3334's B300 evidence, here's H200×4 TP4 with DeepSeek-V4-Flash-0731 (image
The The patch is deployed in our production DSV4 serving cluster (TP4 H200×4). Note: this is arm D (both patches). We have not yet isolated #32467 alone on H200 (arm B) — that test would determine whether P2 is reachable on Hopper or is SM120/B300-specific. @hassellof has proposed a specific set of rows for that cross-check; we plan to run them. |
|
Follow-up B300/TP8 validation: I extended #32467 testing from startup capture to sustained replay, with #33795 kept OFF. Environment: 8x B30Z, TP8, DeepSeek-V4-Pro-DSpark, official v0.5.16-cu130, compact verify, flashinfer_mxfp4. Startup remains clean with #32467 alone: 3/3 fresh servers reached READY and captured all 35/35 decode-graph shapes. Runtime Gate 1 (no SPS, max graph bs=32, max-running-requests=32, c32/n320, 60k/1k): 320/320 completed, 0 timeouts, health 200 before/after, no GPU error signatures, post-run smoke request passed. Runtime sampling showed CUDA Graph replay only (True=134, False=0). Runtime Gate 2' (compact + old SPS, max graph bs=128, c256/n512, 60k/1k, request-timeout=7200s): 512/512 completed, 0 timeouts, health 200 before/after, no illegal-memory / CUDA_ERROR_ILLEGAL_ADDRESS / scheduler-exception / NCCL-watchdog signatures, and the post-run smoke request passed. A second fresh Gate 2' run was manually stopped at 293/512 to release the machine; it had no GPU error signature up to the stop point, so I am not counting it as either pass or failure. Updated conclusion for this B300x8/TP8 reproducer: #32467 alone now has clean evidence for both startup CUDA-Graph capture and sustained replay on the historical strong path. This further supports #32467 as the primary fix for this B300 failure class; #33795 is compatible but was not required for these successful runs. |
Direct kernel-level regression: unpatched produces OOB
|
| input | B | s_min/s_max | actual ragged rows | unpatched | with #32467 |
|---|---|---|---|---|---|
[4] x 96 |
96 | 4/4 | 384 | 0 / 2000 bad | 0 / 10000 bad |
[4] x 72 + [3] x 24 |
96 | 3/4 | 360 | 594 / 2000 bad (29.7%) | 0 / 10000 bad |
[3] x 104 + [2] x 24 |
128 | 2/3 | 360 | 1998 / 2000 bad (99.9%) | 0 / 10000 bad |
Aggregate:
unpatched: 2592 bad plans / 6000 calls
patched: 0 bad plans / 30000 calls
I cleared /root/.cache/tvm-ffi before each arm so both versions were re-JITed from the corresponding c_plan.cuh.
The strongest diagnostic detail is the bad value itself. For both ragged cases the maximum generated ragged_id is exactly:
rmax = 383
= 96 * 4 - 1
= 128 * 3 - 1
= B * s_max - 1
while the real ragged row count is only 360. That is exactly the maximum ID produced by the MTP-uniform formula batch_id * E + j with E=s_max, showing that a genuinely ragged batch was planned as uniform. The uniform control [4] x 96 never produces a bad ID.
So on this B300 reproducer the causal chain is now observable directly at the plan output:
missing __syncthreads()
-> warp scratch init/reduction race
-> ragged input can collapse to the uniform classification
-> uniform-formula ragged_id generation
-> OOB ragged_id
This matches the full-model behavior where the eventual IMA surfaced through different downstream kernels/ranks: those kernels were consumers of an already-invalid plan, not necessarily the producer of the bad address.
Together with the earlier B300 results (#32467-only startup 3/3 READY, sustained replay 320/320 and 512/512 clean), this is much stronger evidence that the barrier in this PR fixes the underlying correctness bug rather than only changing timing.
The harness should also be a good basis for a compact GPU regression test: the two ragged inputs above reproduce quickly without model weights, while the fixed kernel stays clean across 30k calls.
Direct kernel-level regression on B300: unpatched produces OOB plans; this PR eliminates themI added a small single-GPU repro that calls the real DSv4 Environment: The harness intentionally oversizes Results
Aggregate: The first bad ragged plans appear at iteration 0/1, so this is a fast and highly reproducible kernel-level regression rather than a service-level timing probe. Why this directly matches the proposed root causeFor both ragged cases the bad maximum That is the maximum ID generated by the MTP-uniform formula ( The uniform control never produces a bad plan, while the two non-uniform inputs do. This gives a direct reproduction of the failure chain described in this PR: After adding this PR's barrier and clearing the TVM JIT cache so This is consistent with the earlier full-model B300/TP8 result (3/3 startup captures clean with #32467 alone) and the sustained-replay follow-up (320/320 low-tier graph replay + 512/512 on the historical compact/SPS 60k/1k path, with no GPU fault signatures). I have the standalone harness in a form that can be reduced into a 1-GPU pytest regression. I think adding that coverage would be valuable here or as a small follow-up PR. |
Direct kernel-level regression reproducer for #32467I now have a much smaller reproducer that exercises the real GPU JIT path for Environment:
The check is simple: for a ragged input,
Aggregate: I cleared The strongest signature is the bad while the real ragged row count is only 360. That is exactly the maximum ID produced by the uniform-path formula with So this reproducer directly demonstrates the failure chain behind the full-model IMA: This is consistent with the earlier B300/TP8 full-model validation, but is much cheaper and more reproducible: single GPU, tiny inputs, no model weights, and the failure appears within the first few iterations before the fix. I have a Python harness for this and can turn it into a GPU regression pytest if useful for this PR / a follow-up test PR. |
Deterministic single-GPU kernel regression for the missing barrierI added a direct regression harness for Environment:
The harness deliberately oversizes the backing mapping tensors so the race produces inspectable bad plan data instead of immediately crashing the process. The primary check is: I tested three shapes:
Aggregate: For both failing ragged cases, the bad maximum while the actual number of ragged rows is only That is the exact maximum produced by the MTP-uniform indexing formula The uniform control remained clean before and after the patch, so the detector is not simply flagging the normal uniform path. For cache hygiene I cleared the TVM-FFI JIT cache before both arms so the modified One caveat: I did not use the harness's internal This is much cheaper than the whole-model repro (single GPU, tiny tensors, no checkpoint load), so I think it is a good candidate for a GPU regression test. I can send a follow-up PR that converts this harness into a small pytest if maintainers prefer the regression coverage separately from this fix. |
|
Additional producer-level validation following @DarkSharpness's review suggestion to remove the redundant warp-scratch initialization instead of adding another barrier. I ran a same-session A/B/C kernel experiment on B30Z, 10,000 calls per input shape / 30,000 calls per variant:
Inputs: Oracle: This supports the simpler no-initialization formulation under the current fixed launch geometry: each of the 32 warps overwrites its own scratch slot before the existing post-reduction One implementation caveat: if the no-init variant is adopted, I think the fixed 1024-thread / 32-warp assumption should be made explicit with a comment or assertion so a future launch-geometry change cannot silently invalidate the reasoning. Together with the earlier B300 direct harness ( |
Signed-off-by: EanWang211123 <wangyiheng@sangfor.com.cn>
|
Independent B300 runtime regression for the final reviewer-selected no-init implementation. Environment:
The final #32467 implementation from For
Runtime provenance also confirmed that the JIT module used by Stage 1/2 was the no-init variant: Two unrelated integration prerequisites were required before the full DP8 + disaggregation path could run:
I am keeping these separate from the #32467 result. Correctness regression:
Under the natural 60k/1k C64 workload, verify lengths happened to stay uniform, so the ragged producer was not exercised. Using the existing
D-only performance/stability sweep, with debug/assert instrumentation disabled and a fresh engine for each point:
Total: 18,560 / 18,560 successful requests, 0 failures. Post-run health remained 200 for every point, with no preallocation warning or request retraction. This should be read as post-fix runtime no-regression evidence for the final no-init implementation on this B300 matrix, not as a new proof of the producer-race root cause. The producer-level A/B/C regression remains the direct causality evidence. |
|
2026-08-11 follow-up to the earlier final-no-init B300 runtime regression above: I extended the same decode-only topology into an external-concurrency / admission-pressure sweep while keeping
All four arms completed with zero request failures, no observed CUDA illegal-address/device-assert/scheduler faults, no retractions/preallocation warnings, Graph replay retained, and This is not a claim of 1280–2560 simultaneously running requests: MRR remained 1024 and the per-rank running peak remained 128. It is useful as additional sustained-replay/admission-pressure evidence for the final no-init kernel, while true higher in-flight capacity requires raising MRR / graph tiers and re-profiling SPS. The full updated root-cause and runtime/capacity map is in #34297. In particular, the roadmap now records the final #32467 fix as removal of the redundant scratch initialization (not the earlier barrier experiment), and keeps downstream IMA sites classified as consumers/surfacing points. |
|
One review-facing documentation note: the PR diff/HEAD is now the final no-init implementation, but the current PR description still says “Add Current final state: So the current root-cause/fix wording should be: redundant scratch initialization creates a write-write race with per-warp reductions; the final patch removes the redundant initialization and therefore removes the race source. The barrier variant remains useful as an A/B control, not the final implementation. @EanWang211123 if convenient, updating the PR body’s |
|
/tag-and-rerun-ci |
|
cc @weireweire |
…roject#32467) Adopt official fix: remove the warp_min/warp_max init block + __syncthreads() barrier. Each warp writes its own slot via warp::reduce_min/max, making the init redundant and the barrier unnecessary. The existing __syncthreads() after the reduce-write (before warp 0's final cross-warp reduce) ensures all slots are written before any read.
…roject#32467) Signed-off-by: EanWang211123 <wangyiheng@sangfor.com.cn>
… upstream sgl-project#25529) + c_plan barrier (sgl-project#32467) EAGLE draft cuda-graph replay padded batch_size/seq_lens/req_pool_indices/ positions to the captured bucket but left forward_batch.out_cache_loc at the RAW request layout. init_forward_metadata_out_graph then built replay metadata with a padded bs over a raw cache-location buffer — the layout mismatch produces invalid page metadata inside the captured graph and surfaces as an ASYNC illegal memory access at whatever host sync comes first (observed: process_batch_result_decode copy_done.synchronize, alloc_for_decode_prealloc; nondeterministic per TP rank; not reproducible under CUDA_LAUNCH_BLOCKING). This matches the upstream family sgl-project#25512/sgl-project#25529 (same padding signature raw_bs=3 graph_bs=4) and sgl-project#28569 (crash as the running batch shrinks — our mixed load with waves of completing requests triggers exactly the padding boundary; --disable-cuda-graph is immune). Fix: - execute() now swaps forward_batch.out_cache_loc to the zero-filled padded buffers slice (num_tokens * speculative_num_steps) during the padded window, so replay metadata sees the same layout the captured graph does; padded lanes write to reserved cache slot 0. - try/finally restores the raw batch view even when metadata init or replay raises (upstream sgl-project#25529 contract) — a padded leftover poisoned every later consumer of the ForwardBatch. Also ports sgl-project#32467: __syncthreads() between warp-scratch init and the per-warp min/max writes in c_plan.cuh plan_compress_prefill_kernel0 — warp 0 could clobber another warp's slot, misclassifying ragged extend as MTP-uniform and emitting out-of-bounds ragged_id (B300 non-deterministic IMA family sgl-project#33356). Dormant on the GLM path (no callers) but the kernel is shared with DSV4 deployments.
… upstream sgl-project#25529) + c_plan barrier (sgl-project#32467) EAGLE draft cuda-graph replay padded batch_size/seq_lens/req_pool_indices/ positions to the captured bucket but left forward_batch.out_cache_loc at the RAW request layout. init_forward_metadata_out_graph then built replay metadata with a padded bs over a raw cache-location buffer — the layout mismatch produces invalid page metadata inside the captured graph and surfaces as an ASYNC illegal memory access at whatever host sync comes first (observed: process_batch_result_decode copy_done.synchronize, alloc_for_decode_prealloc; nondeterministic per TP rank; not reproducible under CUDA_LAUNCH_BLOCKING). This matches the upstream family sgl-project#25512/sgl-project#25529 (same padding signature raw_bs=3 graph_bs=4) and sgl-project#28569 (crash as the running batch shrinks — our mixed load with waves of completing requests triggers exactly the padding boundary; --disable-cuda-graph is immune). Fix: - execute() now swaps forward_batch.out_cache_loc to the zero-filled padded buffers slice (num_tokens * speculative_num_steps) during the padded window, so replay metadata sees the same layout the captured graph does; padded lanes write to reserved cache slot 0. - try/finally restores the raw batch view even when metadata init or replay raises (upstream sgl-project#25529 contract) — a padded leftover poisoned every later consumer of the ForwardBatch. Also ports sgl-project#32467: __syncthreads() between warp-scratch init and the per-warp min/max writes in c_plan.cuh plan_compress_prefill_kernel0 — warp 0 could clobber another warp's slot, misclassifying ragged extend as MTP-uniform and emitting out-of-bounds ragged_id (B300 non-deterministic IMA family sgl-project#33356). Dormant on the GLM path (no callers) but the kernel is shared with DSV4 deployments.
…roject#32467) Signed-off-by: EanWang211123 <wangyiheng@sangfor.com.cn>
…roject#32467) Signed-off-by: EanWang211123 <wangyiheng@sangfor.com.cn>
8 participants
Motivation
In DSpark compact ragged-verify mode, CUDA graph capture could fail with an illegal memory access in
c128_v2.cuh:506whenmax-running-requestsis large (e.g. 96). The failure was nondeterministic: one TP rank could build a correct plan while another emitted out-of-boundsragged_idvalues.Root cause: in
plan_compress_prefill_kernel0, warp 0 initializes allwarp_min/warp_maxscratch slots while each warp later writes its own slot in Stage B. Without a barrier between init and reduction, warp 0 can clobber another warp's min result. That makes raggedextend_lens(e.g.[4]×72 + [3]×24) look uniform (s_min == s_max), wrongly selecting the MTP fast path and generatingragged_idup tobatch_size × E - 1, past the actual ragged input rows. The write kernel has no bounds check on CUDA, so this surfaces as an IMA.This is a general correctness bug for any non-uniform prefill extend; compact ragged capture just makes it easy to hit.
Modifications
__syncthreads()inc_plan.cuhbetween warp-scratch initialization and per-warp min/max reduction inplan_compress_prefill_kernel0.Accuracy Tests
Speed Tests and Profiling
Setup: DeepSeek-V4-Flash-DSpark, DSPARK block-size 5, TP=2, Marlin MoE,
max-running-requests=96, radix cache disabled. Metric:mean_output_throughput.SGLANG_RAGGED_VERIFY_MODE=compact)compact command:
default command:
Before this fix, compact mode failed during CUDA graph capture at large batch tiers; after the fix, capture completes and serving is stable at bs=96.
Checklist
Review and Merge Process
/tag-and-rerun-ci,/tag-run-ci-label,/rerun-failed-ciCI States
Latest PR Test (Base): ❌ Run #31582596399
Latest PR Test (Extra): ❌ Run #31582595666