Conversation
|
If BCG has been verified in its accuracy, can we enable it by default for GLM5.3 flash model? |
|
BCG should be the default when the resolved stack uses Triton KDA, TRTLLM DSA, and FP8 e4m3 KV now. Chunked prefill and capture ceiling will default to 4096. |
Enable the default after resolving Triton KDA, TRTLLM DSA and FP8 E4M3 KV, while respecting explicit selections and compatibility checks. Default GLM CUDA prefill chunks and breakable capture to 4096 tokens. Handle KDA gate padding, pooled DSA indexing and recurrent-state warmup so GLM text and multimodal prefill can use breakable graphs.
0a4cdc8 to
d4d367a
Compare
415fbc9 to
d4d367a
Compare
|
/rerun-test test/registered/e2e/models/test_glm53_flash_b200.py test/registered/e2e/models/test_glm53_flash_h200.py |
|
Results for 🚀 🚀 |
|
/rerun-test test/registered/e2e/models/test_glm53_flash_b200.py test/registered/e2e/models/test_glm53_flash_h200.py |
|
Results for 🚀 🚀 |
* origin/main: (27 commits) [Simulator] Give the OFFLINE/BLOCKING comparison tolerances real headroom (sgl-project#38732) [Config] msgspec.Struct for the config tier (sgl-project#38753) [AMD] ci: move the miles nightlies from rocm700 to rocm10 (sgl-project#37495) [Config] One writer for the declaration stash; no exception to the write seal (sgl-project#38752) docker(xpu): drop redundant setvars.sh from torch_memory_saver RUN (sgl-project#38665) [XPU][Fix] Pack device-pointer tables as uint64 to avoid 64-bit address overflow (sgl-project#35051) [CI] Temporarily disable GB300 tests (sgl-project#38770) [diffusion] feat: spill large tensors over shared memory like numpy arrays (sgl-project#38656) [diffusion] refactor: refactor utility ownership and document helper placement (sgl-project#38699) [NPU]Support GLM5.2 and FP8 DSA&Indexer kvcache for 950 (sgl-project#38250) [CI] Answer unrecognized slash commands instead of skipping silently (sgl-project#38736) [AMD] Parallelize aiter spec-decode KV index building over token blocks (sgl-project#37659) [DSv4] Integrate TRT-LLM DSv4 Attention for SM100/103 (sgl-project#30805) Add Opt-In for GLM-5.3 Flash breakable prefill CUDA graphs (sgl-project#38522) [CI] Install helion 1.4.0 for the KDA Helion kernel tests (sgl-project#38688) [Rust] Gate health on startup warmup completion (sgl-project#37994) [HiCache] Replace skip_lock_node_ids with a segment lock protocol (sgl-project#36848) feat: add optimized Domino rollout to DFlash V2 (sgl-project#36899) [CI] Add /run-full-ci and /run-extra-ci slash commands (sgl-project#38734) [Model] Support GLM-5.3 Flash NVFP4 loading (sgl-project#38621) ...
…A graph opt-in (hand-merged over local kpool indexer rewrite) (cherry picked from commit 0084030) The forward_cuda/_forward_cuda_impl split was adapted to the local 64-token-page kpool indexer rewrite: the full local body was kept intact under _forward_cuda_impl, with the upstream BCG eager-break wrapper layered on top as forward_cuda and the dual-stream guard added to enable_dual_stream.
The sgl-project#38522 MLA bmm->attention fusion and the 0907 DSA-extend DCP q-gather co-fire on the same BCG extend surface: the fusion materializes q_nope_out in [B,H,D] layout, but all_gather_q_for_mla_decode consumes the bmm-native [H,B,D] written into the DCP combine buffer (use_fused_dcp_q_buf path). Result was a torch.cat dim mismatch (Expected 16 but got 4096) during breakable prefill CG capture warmup on GLM-5.3-Flash (TP4/DCP2, tilelang DSA pair, bf16 KV). Carve the fusion out when dcp_enabled so the fused combine-buffer path runs instead.
…wise) Port and enable breakable prefill CUDA graph under decode context parallelism, including the GLM-5.3-Flash opt-in (sgl-project#38522), NextN draft prefill graphs, the DCP metadata plumbing, and the numerical-correctness fixes found during bring-up: - Port sgl-project#38522 breakable prefill CUDA graph opt-in (hand-merged over the local kpool indexer rewrite) + NextN draft prefill CUDA graph enable - Shared DCP extend metadata builder, persistent DCP buffers, replay-time refresh, and BCG binding of captured segments to those buffers - Set mha_return_lse on DCP extend in prefill CG capture/replay - Preserve head-folded LSE layout / pad LSE to static rows / return folded LSE unchanged in the unified attention op - Narrow-copy shorter eager break outputs into the static bridge buffer - Zero PCG padded tail via narrow, not view - Reshape gathered q in DSA extend absorbed path; reshape non-contiguous attention output copies - Disable BMM-into-attention fusion on DSA-extend under DCP - Infer q head count from q layout in forward_extend under BCG replay: replay re-resolves the 16-head attn_mqa via the layer_id table while q arrives DCP-widened to 32 heads, so layer.tp_q_head_num mis-reshaped q and corrupted all prefill attention output Validated on GLM-5.3-Flash TP4/DCP2 + tilelang DSA: BCG output byte-identical to eager, temp-0 deterministic across repeats, capture collective audit clean under SGLANG_DEBUG_CAPTURE_COLLECTIVE_AUDIT=1.
Motivation
Enable opt-in breakable prefill CUDA graphs (BCG) for GLM-5.3 Flash (
Glm5NextForConditionalGeneration) with Triton KDA, TRTLLM DSA prefill/decode, and FP8 E4M3 KV. Other backend combinations remain eager unless BCG is explicitly selected. Its KDA gates, pooled DSA indexer and recurrent-state operations need capture handling before the architecture can pass the existing default compatibility checks.Task accuracy against ordinary eager
September 9 run:
20260909T170520.859362Z-bcg-latest-main-quality, source8b27ee67f7759e36dc386d04ff2933f5bd7b5e78on mainffe98a4279ba6e42d1f87dc4eeb6edb4887b9ea4. One four-B300 allocation, fresh servers in BCG → eager → eager → BCG order, identical question sets and payloads, fixed two-request waves, no retries or profiling during accuracy measurement.GSM8K uses the full test split with the first five training examples as demonstrations. MMLU samples 1,024 questions with seed 0; the max-reasoning GSM8K subset uses seed 479309393. Output limits are 4,096 tokens for low reasoning and 8,192 for max reasoning.
The validated image is
im-dpE65TM3SbmUUiVLdUKGvC: Transformers 5.16.0.dev0 / Tokenizers 0.23.1 with the GLM image/video processors verified, FlashInfer 0.6.18, DeepGEMM 0.1.7, PyTorch 2.13.0+cu130 and NCCL 2.30.7.Performance
September 7 measurements on release source
b418d1c0387a0f4deb83796c2c0f217feb0be103, using the same four B300 GPUs and BCG → eager → eager → BCG order. Each concurrency point used 250 measured agentic requests per pass, EOS enabled and a 4096-token output ceiling; conditioning was excluded and the prefix cache flushed before each point. These are medians of two passes, not confidence intervals.Interactivity is
1000 / P90 request TPOT_ms, in tokens/s/user. At concurrency 32, P90 TTFT increased from 3.95 s to 4.24 s. These compare serving configurations, including attention dispatch/padding and potentially differing output lengths; they do not isolate graph-launch overhead or measure a live-production speedup.CI States
Latest PR Test (Base): ❌ Run #34417122792
Latest PR Test (Extra): ❌ Run #34417122668
Latest PR Test (AMD ROCm 10): ❌ Run #34417122811