Skip to content

[AMD] [GLM5] Skip DSA decode indexer when kv_len <= index_topk (dense k-only fast path) - #31324

Merged
HaiShaw merged 16 commits into
sgl-project:mainfrom
Jacob0226:jacob/dsa-decode-skip-indexer
Aug 16, 2026
Merged

HaiShaw merged 16 commits into
sgl-project:mainfrom
Jacob0226:jacob/dsa-decode-skip-indexer

Conversation

@Jacob0226

@Jacob0226 Jacob0226 commented Jul 15, 2026

Copy link
Copy Markdown
Contributor

Summary

  • On GLM-5.2 DSA decode, when a request's kv_len <= index_topk the top-k selects all valid positions, so the indexer's logits GEMM + paged_mqa_logits + top-k selection is wasted work. Add a k-only fast path that skips the indexer, stores the K cache, and generates the identity index directly ([0, 1, ..., kv_len-1, -1, ...]), feeding the same sparse-MLA decode attention kernel.
  • CUDA-graph "Design A" dual-graph: capture a dense (k-only) and a sparse (full indexer) decode graph per batch-size bucket and dispatch on max_kv_len vs index_topk at replay. This is correct for mixed lengths — any request with kv_len > index_topk in the batch falls back to the sparse graph.
  • Auto-enabled for DSA models, no env toggle. Since DSA models expose index_topk in their HF config, the dual-graph is turned on automatically whenever index_topk is present (dsa_dual_graph = dsa_index_topk is not None); non-DSA models and other archs are unaffected. Eager decode takes the same fast path per-step when safe. (No SGLANG_DSA_DECODE_* / SGLANG_KONLY_DEBUG environment variables — an earlier revision gated this behind opt-in/debug flags; those were removed per review.)

Changes

File Change
layers/attention/dsa/dsa_indexer.py k-only fast path (_forward_cuda_k_only): skip logits/top-k, store K cache; for MLA generate the identity index; extended to DECODE (eager + graph)
model_executor/runner/decode_cuda_graph_runner.py Design A dual-graph: capture dense (k-only) + sparse decode graphs per bs bucket; host-dispatch on max_kv_len vs index_topk; auto-enabled for DSA models via dsa_index_topk
model_executor/runner_utils/capture_mode.py capture-time DSA variant flag (dense/sparse) read by the indexer skip branch
model_executor/runner/shape_key.py include the DSA decode variant in the cuda-graph shape key

Scope

GLM-5.2 DSA decode on gfx950 (MI355X). Auto-enabled for DSA models; correct for mixed lengths (long context stays on the sparse indexer path). Other archs / hardware unaffected.

Test plan

Accuracy (GSM8K, MI355X TP4, GLM-5.2-MXFP4):

Baseline This PR
GSM8K 0.922 0.941

Within margin of error.

Performance (MI355X TP4, GLM-5.2-MXFP4, docker rocm/sgl-dev:v0.5.15.post1-rocm720-mi35x-20260714, tilelang DSA backend): token throughput per GPU (tok/s/gpu, higher better) and median TPOT (ms, lower better), dense-decode OFF vs ON (this PR). Isolated effect — same build/backend, only this PR's commits added, everything else identical.

Only i1024 / o1024 is shown: this optimization applies when kv_len <= index_topk (2048). At i8192 every decode step has kv_len > 2048, so it stays on the sparse path and is unchanged.

Concurrency TPUT off TPUT PR Δ TPUT TPOT off TPOT PR Δ TPOT
4 71.8 75.5 +5.2% 13.28 12.62 −5.0%
8 126.8 133.2 +5.1% 15.18 14.41 −5.1%
16 196.9 205.1 +4.2% 19.46 18.67 −4.1%
32 312.4 322.8 +3.3% 24.66 23.90 −3.1%
64 468.9 480.3 +2.4% 32.88 32.14 −2.3%
python3 -m sglang.launch_server \
  --model amd/GLM-5.2-MXFP4 \
  --tp 4 \
  --trust-remote-code \
  --tool-call-parser glm47 \
  --reasoning-parser glm45 \
  --watchdog-timeout 1200 \
  --mem-fraction-static 0.85 \
  --kv-cache-dtype fp8_e4m3 \
  --disable-radix-cache \
  --model-loader-extra-config '{"enable_multithread_load": true, "num_threads": 32}' \
  --dsa-prefill-backend tilelang \
  --dsa-decode-backend tilelang \
  --tokenizer-worker-num 8 \
  --enable-aiter-allreduce-fusion

Token throughput per GPU improves ~2.4–5.2% and TPOT drops ~2.3–5.1% at i1k, from skipping the wasted indexer when the sparse top-k would select all positions anyway. The gain is largest at low concurrency and converges to ~2.4% at conc 64. (Δ is the isolated dense-decode effect — same build/backend, feature off vs on — not a full-stack-vs-stock comparison.)


CI States

Latest PR Test (Base): ⏳ Run #31978978932
Latest PR Test (Extra): ❌ Run #31978978765

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request introduces a dense-decode optimization for DSA models, allowing the indexer's logits computation to be skipped (k-only path) during decode when the sequence length is within the top-k threshold. This is supported in both eager mode and via a dual-graph CUDA graph capture mechanism. The review feedback highlights opportunities to improve robustness by using the central envs registry instead of direct os.environ.get calls (ensuring proper boolean parsing) and adding .numel() > 0 checks to prevent RuntimeErrors when calling .max() on empty sequence length tensors.

Important

The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.

Comment on lines +577 to +582
# M2a "static" fallback (no dual-variant capture active): enable
# k-only only when the deployment guarantees every request's
# kv_len<=index_topk (e.g. i1k/o1k). WRONG for mixed >2K traffic.
import os

return os.environ.get("SGLANG_DSA_DECODE_DENSE_GRAPH", "0") == "1"

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

Using os.environ.get directly bypasses the envs configuration registry. More importantly, os.environ.get(..., "0") == "1" only evaluates to True if the environment variable is exactly "1". If a user sets SGLANG_DSA_DECODE_DENSE_GRAPH=True or SGLANG_DSA_DECODE_DENSE_GRAPH=true, this check will evaluate to False. Using envs.SGLANG_DSA_DECODE_DENSE_GRAPH.get() correctly utilizes the central environment variable registry and its boolean parsing logic.

Suggested change
# M2a "static" fallback (no dual-variant capture active): enable
# k-only only when the deployment guarantees every request's
# kv_len<=index_topk (e.g. i1k/o1k). WRONG for mixed >2K traffic.
import os
return os.environ.get("SGLANG_DSA_DECODE_DENSE_GRAPH", "0") == "1"
# M2a "static" fallback (no dual-variant capture active): enable
# k-only only when the deployment guarantees every request's
# kv_len<=index_topk (e.g. i1k/o1k). WRONG for mixed >2K traffic.
return envs.SGLANG_DSA_DECODE_DENSE_GRAPH.get()

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done — switched to envs.SGLANG_DSA_DECODE_DENSE_GRAPH.get() in commit 474270e so it goes through the EnvBool parser (accepts true/1/yes).

Comment on lines +1360 to +1367
if forward_batch.forward_mode.is_decode_or_idle() and layer_id == 0:
import os

if os.environ.get("SGLANG_KONLY_DEBUG", "0") == "1":
logger.info(
"[KONLY] decode k-only fired (skip-indexer): num_tok=%d",
x_meta.shape[0],
)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

Using os.environ.get directly bypasses the envs configuration registry. More importantly, os.environ.get(..., "0") == "1" only evaluates to True if the environment variable is exactly "1". If a user sets SGLANG_KONLY_DEBUG=True or SGLANG_KONLY_DEBUG=true, this check will evaluate to False. Using envs.SGLANG_KONLY_DEBUG.get() correctly utilizes the central environment variable registry and its boolean parsing logic.

        if forward_batch.forward_mode.is_decode_or_idle() and layer_id == 0:
            if envs.SGLANG_KONLY_DEBUG.get():
                logger.info(
                    "[KONLY] decode k-only fired (skip-indexer): num_tok=%d",
                    x_meta.shape[0],
                )

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done — switched to envs.SGLANG_KONLY_DEBUG.get() in commit 474270e, same EnvBool parsing.

Comment thread python/sglang/srt/layers/attention/dsa/dsa_indexer.py
Comment on lines 585 to 591
if fb.seq_lens_cpu is not None:
max_kv_len = int(fb.seq_lens_cpu.max().item())
elif fb.seq_lens is not None:
max_kv_len = int(fb.seq_lens.max().item())
else:
return False
return max_kv_len <= self.index_topk

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

If fb.seq_lens_cpu or fb.seq_lens is empty (i.e., has 0 elements), calling .max() on it will raise a RuntimeError. Adding a .numel() > 0 check ensures robustness against empty batches.

Suggested change
if fb.seq_lens_cpu is not None:
max_kv_len = int(fb.seq_lens_cpu.max().item())
elif fb.seq_lens is not None:
max_kv_len = int(fb.seq_lens.max().item())
else:
return False
return max_kv_len <= self.index_topk
if fb.seq_lens_cpu is not None and fb.seq_lens_cpu.numel() > 0:
max_kv_len = int(fb.seq_lens_cpu.max().item())
elif fb.seq_lens is not None and fb.seq_lens.numel() > 0:
max_kv_len = int(fb.seq_lens.max().item())
else:
return False
return max_kv_len <= self.index_topk

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 27bd826 — added .numel() > 0 checks to both the seq_lens_cpu and seq_lens branches before .max(), consistent with the existing empty-batch guard elsewhere in this file.

…ger)

When kv_len <= index_topk the DSA top-k selects all valid positions, so the
indexer logits GEMM + paged_mqa_logits + top-k are pure waste. Route decode
through the existing _forward_cuda_k_only path (dummy-logits topk_transform
-> correct physical page-slot indices), instead of the spike's arange()
identity synthesis which was wrong under SGLANG_DSA_FUSE_TOPK (indices are
physical slots, not logical positions) -> GSM8K 0.02.

- _should_skip_logits_computation: allow is_decode_or_idle() with the same
  max_kv_len <= index_topk test; gated off during cuda-graph capture (the
  captured branch would be frozen and replay wrong for kv_len>index_topk).
- _forward_cuda_k_only: relax the extend-only assert to also accept decode.

Eager-mode only for now (Milestone 1: correctness). Cuda-graph dense/sparse
variant + can_run_graph gating is Milestone 2.
…a-graph decode k-only (kv_len<=index_topk deployments)
…t dispatch

Capture two decode graph variants per bs bucket -- "dense" (k-only,
skip-indexer) and "sparse" (full indexer) -- and select which to replay
on the host from max_kv_len vs index_topk. Correct for mixed/arbitrary
context length under cuda graph (>2K routes to sparse, <=2K to dense),
never falling back to eager.

- capture_mode.py: add _capture_dsa_variant signal (mirrors lora variant)
- dsa_indexer.py: _should_skip_logits_computation capture branch honors the
  variant (dense->skip, sparse->full); M2a static env kept as fallback
- shape_key.py: add dsa_variant to ShapeKey (composes with lora variant_label)
- decode_cuda_graph_runner.py: capture dense+sparse per bs; _resolve_dsa_variant
  dispatches on host seq_lens_cpu (no d2h); env-gated dispatch debug log
- environ.py: register SGLANG_DSA_DECODE_DUAL_GRAPH / _DENSE_GRAPH / KONLY_DEBUG

Validated (GLM-5.2-MXFP4 TP4 MI355X, index_topk=2048): GSM8K under graph
0.950; long >2K needle recall correct; i1k dense -3~5% TPOT; i8k no
regression; dual capture ~55s (~2x), no OOM.
@Jacob0226
Jacob0226 force-pushed the jacob/dsa-decode-skip-indexer branch from d8c3c81 to 0c41aca Compare July 16, 2026 08:09
SGLANG_DSA_DECODE_DENSE_GRAPH and SGLANG_KONLY_DEBUG are declared as
EnvBool (accepts true/1/yes/y), but dsa_indexer.py read them with
os.environ.get(...) == "1", so common forms like
SGLANG_DSA_DECODE_DENSE_GRAPH=true silently no-op'd. Use envs.<FLAG>.get()
for both so behavior matches the EnvBool declaration and the rest of the
codebase's env conventions.

Co-authored-by: Cursor <cursoragent@cursor.com>
Jacob0226 and others added 3 commits July 19, 2026 21:50
…ner (--mtp)

EAGLEDraftCudaGraphRunner subclasses DecodeCudaGraphRunner and reuses its
capture(), but does not run DecodeCudaGraphRunner.__init__ (so it never sets
dsa_dual_graph) and overrides capture_one_shape with a signature that has no
dsa_variant. Under --mtp this raised
  AttributeError: 'EAGLEDraftCudaGraphRunner' object has no attribute 'dsa_dual_graph'
at cuda-graph capture. Guard the dual-graph checks with getattr(..., False) and,
for the None variant, call capture_one_shape without the extra dsa_variant arg so
the narrower draft override still works.

Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

# Conflicts:
#	python/sglang/srt/model_executor/runner/decode_cuda_graph_runner.py
Add .numel() > 0 checks before calling .max() on seq_lens_cpu / seq_lens
in _should_skip_logits_computation, so an empty batch cannot raise a
RuntimeError. Consistent with the existing numel guard elsewhere in this
file. Addresses review feedback on the extend and eager-decode gates.

@clintg6 clintg6 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@Jacob0226 Tested on MI355X with GLM-5.2-MXFP4 (TP8, 1024 ISL / 1024 OSL). The dual-graph dispatch works correctly. Benchmark results showed around 3% uplift in Total token throughput. With that said, I have a number of recommendations to improve the PR.

Remove SGLANG_KONLY_DEBUG

This appears to be a leftover debug flag. It's only used to gate a single logger.info call in _forward_cuda_k_only, so it shouldn't ship in production. Remove it from environ.py and the associated logging block in dsa_indexer.py.

Remove SGLANG_DSA_DECODE_DUAL_GRAPH and auto-enable for DSA models

Since DSA models already expose index_topk in their HF config, users shouldn't need to enable this manually. Auto-detect DSA models instead:

# Instead of:
self.dsa_dual_graph = envs.SGLANG_DSA_DECODE_DUAL_GRAPH.get()

# Use:
self.dsa_dual_graph = self.dsa_index_topk is not None

Remove SGLANG_DSA_DECODE_DENSE_GRAPH

This forces dense decode by unconditionally bypassing the indexer, making it another debug/testing override and is unnecessary.

@Jacob0226

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

@Jacob0226

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

@Jacob0226

Copy link
Copy Markdown
Contributor Author

@amd-bot ci-status

@amd-bot

amd-bot commented Aug 13, 2026

Copy link
Copy Markdown

@Jacob0226

CI Status for PR #31324

Merge verdict:Do not merge on green alone. PR CI is incomplete (AMD downstream fast-fail skipped) and this PR's core changed code is not exercised by any PR-CI test. All 4 executed failures are pre-existing/infra issues unrelated to this DSA decode change — but "0 related failures" here does not mean the PR is verified. The new k-only / dual-graph decode path is AMD-only + DSA-model-gated, and nothing in PR CI loads a DSA model on AMD.

Caution

Two coverage gaps, both blocking verification:

  1. Changed code is unreachable by PR CI. The feature activates only when is_hip() and is_deepseek_dsa(hf_config) (decode_cuda_graph_runner.py) — i.e. AMD GPU + a DSA model (DeepSeek-V3.2 / GLM-5.2 with index_topk). All DSA/GLM-5 end-to-end tests live under test/manual/ (README: "Non-CI tests") and require B200/GB300/8-GPU. No registered PR-CI suite loads a DSA model on AMD. Green proves only that the modules import and the non-DSA graph path is unbroken.
  2. AMD pipeline did not finish. wait-for-stage-b-amd failed (cascade from an unrelated test_tracing failure in stage-b-test-1-gpu-small-amd (9)), which skipped 11 downstream AMD jobs — including stage-c-test-4-gpu-amd, stage-c-test-large-8-gpu-amd, and stage-c-test-large-8-gpu-amd-mi35x, the multi-GPU suites where a large DSA model could plausibly run. AMD signal for this PR is not full.

Before merge: run a DSA model (GLM-5.2 / DeepSeek-V3.2) decode on AMD MI300 to actually exercise _forward_cuda_k_only + the dense/sparse dual-graph dispatch (both the k-only and the mixed kv_len > index_topk fallback). Re-run AMD CI (rebase past the test_tracing fix, or apply bypass-fastfail) so stage-c completes.

Changed files: dsa/dsa_indexer.py (+67/−6), runner/decode_cuda_graph_runner.py (+91/−10), runner/shape_key.py (+6/−3), runner_utils/capture_mode.py (+18/−0)

Executed CI failure attribution: AMD: 1 executed failure (0 related) · NPU: 2 (0 related) · MLX: 1 (0 related). No jobs pending. (11 AMD downstream jobs skipped by fast-fail — counted as a completeness gap above, not as failures.)

AMD Executed Failures

Job Test File Test Function Error Related? Why
stage-b-test-1-gpu-small-amd (9) test/registered/observability/test_tracing.py TestTraceServerAsync.setUpClass AttributeError: module 'sglang.srt.observability.trace' has no attribute 'global_trace_level' → server 500 on warmup 🟢 unlikely Observability/tracing module; PR touches only DSA indexer + CUDA-graph runner. Pre-existing on main. This is what tripped wait-for-stage-b-amd and cascaded the AMD skips.

Other Executed Failures

Job Test File Test Function Error Related? Why
base-c-test-perf-8-npu-a3 test/registered/npu/performance/minimax_m2_5/test_npu_minimax_m2_5_w8a8_4p_in64k_out1k_prefix90_50ms.py perf assertion AssertionError: 52.31 not less than or equal to 51.0 🟢 unlikely NPU MiniMax-M2.5 latency threshold miss; unrelated backend + model.
multimodal-gen-test-2-npu-a3 sglang/multimodal_gen/test/server/ascend/test_server_2_npu.py TestDiffusionServerTwoNpu.test_diffusion_generation[*] RuntimeError: Server exited early (code 1) — scheduler dead, Exit code: -9 (SIGKILL/OOM) 🟢 unlikely NPU diffusion (multimodal_gen) server OOM; separate code path from LLM DSA decode.
stage-a-unit-test-mlx test/registered/unit/hardware_backend/mlx/test_metal_profiler.py Metal capture test AssertionError: Failed to start MLX Metal capture: Capture layer is not inserted 🟢 unlikely Apple-Silicon MLX profiler infra (needs MTL_CAPTURE_ENABLED=1); unrelated backend.

Cascade/aggregator failures collapsed (not independent): pr-test-{npu,amd,mlx,extra,amd-extra}-finish, wait-for-stage-b-amd, call-gate / pr-gate (×2).

Details / what to do before merge

  • Verify the changed code for real — the single most important gap. Launch a DSA model decode on AMD MI300 and confirm: (a) k-only fast path when kv_len <= index_topk, (b) correct sparse fallback for a mixed batch with some kv_len > index_topk, (c) the dense/sparse dual-graph capture + _resolve_dsa_variant replay dispatch. The registered unit tests that import the changed modules (test_decode_cuda_graph_war_fence.py, test_hidden_state_graph_recapture.py) passed but do not set dsa_dual_graph (no DSA model, not HIP), so they exercise only the dsa_variant=None default paths.
  • Get full AMD signal — the test_tracing global_trace_level failure is pre-existing (In-flight fix: ❌ none found in open PRs) and is what skipped stage-c 4/8-GPU AMD. Rebase past a tracing fix or apply bypass-fastfail (uses more CI, use sparingly) and re-run so downstream AMD jobs actually execute.
  • NPU + MLX failures need no action from this PR's author (unrelated backends/infra), but the NPU diffusion OOM and MiniMax perf miss are worth flagging to the NPU owners separately.

Generated by amd-bot using Claude Code CLI

@Jacob0226

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

@Jacob0226

Copy link
Copy Markdown
Contributor Author

@amd-bot ci-status

@amd-bot

amd-bot commented Aug 14, 2026

Copy link
Copy Markdown

@Jacob0226

CI Status for PR #31324

Merge verdict: No blocking PR-caused failures — all 3 executed test failures are in code paths this PR does not touch (AMD allreduce-fusion mock, NPU MiniMax perf, NVIDIA NIXL RDMA infra). But do not merge on "green" alone: this PR's entire value — the DSA decode skip-indexer fast path for GLM5/DeepSeek-V3.2 on AMD — is exercised only by nightly=True suites, so PR CI never ran the changed code.

Caution

This PR's changed code is not exercised by any PR-CI test. The DSA decode-indexer fast path in dsa_indexer.py / decode_cuda_graph_runner.py is only covered by AMD accuracy suites registered nightly=True (nightly-amd-accuracy-8-gpu-glm5, nightly-amd-accuracy-8-gpu-deepseek-v32, nightly-amd-8-gpu-mi35x-glm5-mxfp4, nightly-amd-accuracy-8-gpu-glm51). None of these run on PR CI. A green run does not verify this change. Before merging, run a GLM5 / DeepSeek-V3.2 decode eval on AMD (e.g. test/registered/amd/accuracy/mi30x/test_glm5_eval_amd.py or test_deepseek_v32_eval_amd.py with CUDA-graph decode enabled) and confirm the kv_len <= index_topk dense k-only fast path produces correct output.

Changed files: dsa/dsa_indexer.py (+67/-6), runner/decode_cuda_graph_runner.py (+91/-10), runner/shape_key.py (+6/-3), runner_utils/capture_mode.py (+18/-0)

Executed CI failure attribution: AMD: 1 failure (0 related) · Others: 2 root failures (0 related) — NPU perf2/4/16 are fast-fail cascades of NPU perf8; B200 shard cancelled (infra); all *-finish / pr-gate jobs are aggregators reflecting the above.

AMD Executed Failures

Job Test File Test Function Error Related? Why
stage-c-test-large-8-gpu-amd-rocm720 (mi300-8gpu, 3) test/registered/ops/test_aiter_allreduce_fusion_amd.py TestAiterAllreduceFusionGate.test_dense_tp_fuses AttributeError: 'types.SimpleNamespace' object has no attribute 'moe_ep_size' 🟢 Failure is in the aiter allreduce-fusion gate; the test's SimpleNamespace mock lacks moe_ep_size. PR touches DSA decode indexer / CUDA-graph capture only — no allreduce, MoE, or server-args code. Test-mock/prod drift unrelated to this PR.

Other Executed Failures

Job Test File Test Function Error Related? Why
base-c-test-perf-8-npu-a3 test/registered/npu/performance/minimax_m2_5/test_npu_minimax_m2_5_w8a8_4p_in64k_out1k_prefix90_50ms.py test_npu_minimax_m2_5_w8a8_4p_in64k_out1k_prefix90_50ms AssertionError: tpot 51.86 not <= 51.0 (baseline 50.0) 🟢 NPU (Ascend) backend, MiniMax-M2.5 W8A8 perf threshold miss by 1.7%. Different backend & model; PR is AMD/GLM DSA decode. Marginal perf flake. NPU perf2/4/16 fast-fail-cascaded off this root cause.
base-c-test-8-gpu-h20 (1) test/registered/disaggregation/test_disaggregation_nixl.py TestDisaggregationNixlAccuracy.setUpClass UCX ERROR mlx5dv_devx_general_cmd ... Remote I/O errorNIXL_ERR_BACKEND → server exit -9 🟢 NIXL/UCX InfiniBand RDMA registration failure on the h20 runner — a known-flaky infra condition (CLAUDE.md notes h20 dirty-GPU state is filtered from fast-fail). From an earlier attempt (2026-08-13). Not code-related.

base-c-test-4-gpu-b200 (2) was cancelled (A task was canceled during action download) — infra, not a test failure.

Details / what to do before merge

  • Coverage (the real action item): This is a decode-time behavior change (skip the DSA indexer when kv_len <= index_topk). It only triggers when serving a DSA model (GLM5 / DeepSeek-V3.2) in decode with CUDA graph — a path with zero PR-CI coverage since all such suites are nightly=True. Trigger the relevant nightly AMD accuracy suite (or run test_glm5_eval_amd.py / test_deepseek_v32_eval_amd.py manually on MI300/MI325) before merging; confirm both the fast-path (kv_len <= index_topk) and the normal indexer path produce matching accuracy.
  • AMD allreduce failure: not caused by this PR, but it is a real red on rocm720 stage-c — worth a separate check by the AMD team (looks like an aiter allreduce-fusion gate test whose mock is missing moe_ep_size; may already be a known/tracked break).
  • NPU / h20 / b200 failures: unrelated infra/perf flakes on other vendors; no action needed for this PR.

Generated by amd-bot using Claude Code CLI

@HaiShaw

HaiShaw commented Aug 16, 2026

Copy link
Copy Markdown
Collaborator

CI green, final update to comment 903fe7d

@HaiShaw
HaiShaw merged commit f7cb328 into sgl-project:main Aug 16, 2026
66 of 98 checks passed
danielchristiancazares pushed a commit to danielchristiancazares/sglang that referenced this pull request Aug 17, 2026
… k-only fast path) (sgl-project#31324)

Co-authored-by: Thomas Wang <thomawan@amd.com>
Co-authored-by: HaiShaw <hixiao@gmail.com>
jybsuper added a commit to jybsuper/sglang that referenced this pull request Aug 24, 2026
A grouped-GEMM MoE LoRA engine selected by --moe-runner-backend lora.
Adapters ride the base expert GEMMs instead of the virtual-experts
fused path:

- Providers own the base stages (prepare/gateup/activation/down/
  finalize) over two row domains: masked [E, m_max, *] slabs for decode
  and flat aligned segments for prefill. Vendors: CuTeDSL (default,
  SM90/SM100 dual-stage packed-schedule kernels), DeepGEMM, and a
  route-major Triton arm.
- LoRA deltas fuse into the stage seams: gate_up delta lands pre-SwiGLU
  in the activation kernel, down-B rides the base GEMM, finalize merges
  routed scaling; per-expert and shared-outer adapter layouts.
- Plan/tile selection from JSON config tables (configs/, overridable
  via SGLANG_LORA_MOE_CONFIG_DIR) keyed by architecture, phase, and
  rank band; construction-time validation, no silent fallbacks.
- Graph-stable workspaces with eager/capture buffers and side-stream
  overlap; spec-aware batch metadata via get_batch_token_counts.
- Unit + registered suites (282 tests) and the lora_moe benchmark
  harness.

[AMD][CI] Add the Qwen3.8 MXFP4 MI35x nightly (#35383)
[Mamba] fix mamba index h unexpected assertion for dcp (#36005)
[AMD] Update amd deepseek v4 cookbook 0822 (#35854)
[diffusion] feat: support loading native diffusers miniMax h3 components (#36067)
[diffusion] feat: support hybrid conditioning for minimax h3 (#36080)
[npu] Kill evalscope session by process group and fix report score parsing (#35988)

Co-authored-by: sglang-npu-bot <sglangnpu@163.com>
[diffusion] feat: support compact qwen3-vl conditioning for minimax h3 (#36076)
[NPU] [DOC] Add Ascend NPU (A3) recipe to the Kimi-K3 cookbook (#35508)
[diffusion] CI: guard the anonymous-host budget alongside peak VRAM (#36051)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
[diffusion] feat: automatically infer comfy fp8 activation scaling (#36060)
[NVIDIA] Fix SM107 MXFP8 activation prep (#35405)

Signed-off-by: Sahithi Chigurupati <chigurupati.sahithi@gmail.com>
Co-authored-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>
[Fix] lfm2 detector: recover tool calls dropped by common model-outpu… (#34237)

Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: Xinyuan Tong <xinyuantong.cs@gmail.com>
[diffusion] UX: clean up startup and offload logs (#36034)
config: publish before the launcher reads effective configuration (#35910)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
config: pin two orderings resolution relies on (#35909)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
config: borrowed-record reads follow the config bags (#35908)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
config: constructing a config no longer resolves it (#35907)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
config: project the config bags from the resolution result (#35906)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
config: record resolution writes in a declaration stash (#35905)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
config: a defensive publish must not re-project over a live process (#35904)
[diffusion] CI: let the 5090 consumer case runs two warm requests on the full recipe (#36032)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
[AMD][DSV4] perf: use full 1024-thread block for indexer top-k on ROCm (#36004)
[AMD] Add Radix-4 MoE top-k router kernel for Kimi-K3 routing (#34490)
refactor(disagg): move _is_watermark_ready into StagingManagerMixin (#36030)
refactor(disagg): dedupe mooncake failure_exception into a mixin (#36031)
[NPU] [DOC] Refresh supported features and models on Ascend NPU (#35836)
refactor(disagg): register SGLANG_ENCODER_MM_LOAD_WORKERS in Envs (#36006)
[diffusion] feat: release a layerwise component's non-layer weights between uses (#35734)
[diffusion] feat: support loading serialized comfy convrot int8 native encoders (#36023)
[diffusion] fix: stabilize ltx-2.3 two-stage cold requests (#35997)
[diffusion] chore: re-home decode-dtype vae weights to a file-backed store (#35986)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
[diffusion] feat: support serialized comfy convrot int8 dits (#35994)
[Kimi-K3] Fix "wrong grids" crash in DP-sharded vision preprocessing (#35305)

Co-authored-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>
[AMD][Spec] Fix aiter GQA packing + split-KV routing in NEXTN spec attention (verify & draft_extend) (#30105)
refactor(disagg): hoist staging helper imports out of the bootstrap loops (#35980)
[diffusion] docs: add tuning guide for h3 on consumer-level gpu (#35816)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[NPU] [Diffusion] Fix critical Ascend NPU Diffusion regression/bugs & restore 2-NPU CI testcase (#34855)

Co-authored-by: Elizaveta Martirosian <elizabet3000@mail.ru>
Co-authored-by: Arseniy Mironov <98156294+Napkin-AI@users.noreply.github.com>
Co-authored-by: Alexandr <117110413+Allor-maker@users.noreply.github.com>
Co-authored-by: P_Alex_Tr <aleksandr.smyshlaev@yandex.ru>
[diffusion] comfyui: add a minimax-h3 node and a generic extra-fields passthrough (#35352)
[diffusion] feat: support single-file component weight overrides (#35979)
[diffusion] fix: fix hunyuan3d stale extension lock hangs (#35989)
[AMD] DeepSeek-V4: add aiter fused mHC post+pre with cross-layer boundary dispatch (#32577)

Co-authored-by: 1am9trash <1am9trash@gmail.com>
Co-authored-by: HAI <hixiao@gmail.com>
[diffusion] chore: let the auto policy select h3's dit for layerwise offload (#35812)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[diffusion] optimization: keep vae decoder weights in their decode dtype from load (#35967)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
fix(test): unbreak test_kv_transfer_replica_metric after #35950 (#35974)
[diffusion] feature: use the directory for the vae mapping gate (#35946)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[diffusion] feat: admit compatible quantized native encoders (#35962)
[diffusion] optimization: transfer mapped layers through a courier thread (#35882)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[FEAT] Weight Daemon abstraction (#33279)

Co-authored-by: liusy58 <liusy58@smail.nju.edu.cn>
refactor(disagg): drop dead placeholder overrides in Common KV sender/receiver (#35950)
refactor(disagg): hoist duplicated _handle_staging_req into a mixin (#35948)
[diffusion] feat: load serialized bnb4 components with transformers (#35945)
[VLM] Split Pixtral multi-image features before the CUDA IPC wrap (#35463)
Make draft attention backends extensible (#35932)

Co-authored-by: Yichao Fu <yichaofu@meta.com>
[Model] Complete dots.note.omni support with native encoders, video preprocessing, and MTP decoding (#33829)

Co-authored-by: miraclezqc <dysania@pku.edu.cn>
[HiCache] Clamp tombstoned SWA locs in UnifiedSWAKVPool translation (#35933)
[docs] Re-measure the Qwen3.8-27B RTX 5090, RTX PRO 6000 and DGX Spark grids on 1cf2b8c (#35825)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
[diffusion] Enable SANA-Video breakable CUDA graphs (#35729)
[diffusion] Fuse SANA-Video interleaved RoPE (#35695)
[CI] Re-enable B300 jobs (#35607)

Co-authored-by: Mohammad Angkad <mohammad.angkad@radixark.ai>
[diffusion] feat: resolve hub component subfolders (#35939)
fix(disagg): PD transfer-failure injection was silently inert (#35890)
Fix buffer-mode HiCache load-back ownership races; add optional prefetch anchor lock (#35769)

Signed-off-by: Zhiqiang Xie <zqx@meta.com>
Add sampling observer auxiliary output hooks (#35747)

Co-authored-by: Alec Solder <alecs@fb.com>
Support CPU offload for mxfp8 KV cache (#35888)
[MLX] Upgrade to Torch 2.13/MLX 0.32+ and redesign the Torch-MLX tensor bridge (#32984)

Co-authored-by: Alex Nails <alex.nails@radixark.ai>
[DeepSeek V4] Add W4A4 MegaMoE server flag (#35918)
[diffusion] refactor: hand out pinned host memory per layer (#35867)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[diffusion] feat: reject unsupported quantized component checkpoints (#35873)
[diffusion] feat: keep a cpu-started vae weights on the checkpoint mapping (#35862)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[Fix] Read the granite sinks dtype from the exec bag, not the legacy global shim (#35921)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
[DSA] Route the ragged prefill top-k to the v2 kernel (#35175)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
[DeepSeek V4] Default FP4 checkpoints to FlashInfer MXFP4 MoE (#35919)
Rainj me/rust server refactor2 (#35239)
[Docs] Add --prerelease=allow to cookbook uv install commands (#35920)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
docs: add website link to README header (#35210)
perf: overlap Qwen shared expert with DeepEP routed experts (#34938)
[Spec][LoRA] Support multi-adapter LoRA with EAGLE/NEXTN/DFLASH/DSPARK speculative decoding (#34337)
[Runtime] Don't override CUDA_MODULE_LOADING (#35711)

Co-authored-by: TRAE CLI <traecli@bytedance.com>
chore: bump docs install version to 0.5.18 (#35911)

Co-authored-by: sglang-bot <sglang-bot@users.noreply.github.com>
docs: add DSPARK speculative decoding option to Ling-3.0-flash cookbook (#35861)
refactor(disagg): extract _all_reduce_polls helper (#35886)
fix(grpc): derive choice count before normalization (#35778)

Signed-off-by: Connor Carpenter <connorc@nvidia.com>
[Fix] Pass Anthropic thinking history as reasoning_content for custom chat encoders (#35480)

Co-authored-by: Mohammad Angkad <mohammad.angkad@radixark.ai>
Add SGLang Granite SWA support via existing Granite models (#35794)

Signed-off-by: Davis Wertheimer <davis.wertheimer@ibm.com>
[mem_cache] docs: add a layer map and placement rules (#35643)

Co-authored-by: ispobock <ispobaoke@gmail.com>
Restructure mem_cache auto-labels by layer (#25122)

Co-authored-by: ispobock <ispobaoke@gmail.com>
[diffusion] feat: support loading peft lora (#35868)
[diffusion] fix: fall back to a component's default attention backend (#35796)

Co-authored-by: Mick <mickjagger19@icloud.com>
[diffusion] fix: do not warn that the recommended short edge is unverified (#35745)
refactor(disagg): collapse duplicated branches in get_kv_class (#35847)
[diffusion] fix: fix quantized qkv scales and missing-param policy for minimax-h3 (#35740)
refactor(disagg): remove unreferenced dead code (#35838)
[diffusion] refactor: resolve lora weight sources deterministically (#35774)
[diffusion] fix: stop the mapped-weight store from holding the parameter itself (#35813)
refactor(disagg): remove dead get_embedding_port (#35844)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
[diffusion] Accelerate SANA-Video linear attention in quality=high (#35728)
[diffusion] Enable LongCat breakable CUDA graphs (#35724)
refactor(disagg): remove dead build_and_send_encode_request (#35843)
Fix overlap prebuilt row reuse race (#35748)
[Fix] Clear full-to-SWA mapping with `index_fill_` to avoid a blocking H2D copy (#35773)
[AMD] fix(rocm): support flydsl 0.3.0 in the FlyDSL fused norm kernel (#34536)

Co-authored-by: Bingxu Chen <195740905+bingxche@users.noreply.github.com>
Co-authored-by: thomawan <thomawan@amd.com>
[mem_cache][9/N] refactor: move DSAIndexerPoolHost to pool_host.dsa (#35306)
[Refactor] New EPD (#30398)

Co-authored-by: Yuang Chen <1131578721@qq.com>
Co-authored-by: Yuang Chen <cya539102@antgroup.com>
Co-authored-by: ZhengWG <zwg0606@gmail.com>
[AMD] Update ROCm AITER pin to c16d44b (#35810)
add py env activate in xpu kernel release workflow (#35726)
[AMD] DSv4: fuse the qk-norm-rope pair on the MTP target-verify path (#34973)

Co-authored-by: HAI <hixiao@gmail.com>
Co-authored-by: Thomas Wang <thomawan@amd.com>
[AMD][CI] Fix ROCm 7.0's dead apt index fail the MORI dependency install (#35764)

Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: quitenode <quitenode@users.noreply.github.com>
Support mxfp8 KV cache in PD transfer (#35718)
[AMD] CI: cut two setup cycles from the AMD multimodal-gen lanes (#34483)

Co-authored-by: quitenode <quitenode@users.noreply.github.com>
[docs] Retune the Qwen3.8-27B RTX 5090 DFLASH2 cells against 1cf2b8c (#35786)
[AMD] Retry transient network failures in ROCm Dockerfile curl fetches (#35654)
[AMD] Improve K3 dspark draft attn kernel perf (#35499)
[AMD] DeepSeek-V4 MI355X: eliminate bpreshuffle fp8-scale copies at producer sites (MoE down, MLA o_proj bmm) (#33166)

Co-authored-by: kk <43161300+kkHuang-amd@users.noreply.github.com>
Co-authored-by: Thomas Wang <thomawan@amd.com>
[diffusion] feat: allow offloaded weights stay on the checkpoint mapping (#35701)
Using unified radix tree by default for all case (#35081)
[doc] standardize diffusion cookbook model pages (#34247)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[diffusion] Fuse LTX-2.5 decoder 3D RoPE (#35698)
[P/D disagg] Decode-side radix cache for SWA hybrid models (unified radix tree) (#27770)

Co-authored-by: Shangming Cai <csmthu@gmail.com>
[CI] Temporarily disable B300 jobs (#35627)
[diffusion] chore: read the cgroup this process is actually in (#35707)
[Sampling] Restore finite top-k requirement for sampling masks (#35205)
[docs] Point the Qwen3.8-27B DFLASH2 note back at the rolling dev image tag (#35767)
fix(kernel) Fix Helion small-token prefill bug (#35197)

Co-authored-by: Ethan Che <eche@meta.com>
[docs] fix note formatting in sglang-d documentation (#35761)
[docs] Tell Qwen3.8-27B DFLASH2 users to build from main (#35753)
Fix _GenerationStreamAccumulator logprob_end off-by-one under retract (#26510)

Co-authored-by: Qiaolin Yu <liin1211@outlook.com>
[AMD] [sgl-kernel] Bypass caches for peer traffic in ROCm custom all-reduce (#32832)

Co-authored-by: Hubert Lu <Hubert.Lu@amd.com>
fix(openai): avoid duplicate routed expert in response when `return_meta_info = True` (#35323)
Add CI permissions for four contributors (#35600)
[CI] Gate `/rerun-test` on commenter trust and remove `/rerun-stage` (#35750)
[Docker] Defer CUDA 13 NCCL override until after dependency resolution (#35756)
feat: make mm_inputs msgpack-native (#29656)

Co-authored-by: Alex Nails <alex.nails@radixark.ai>
[misc] Trim restating comments and docstrings in srt/managers (#35622)
[Kimi K3] Select FlashInfer MXFP4 for SM107 auto MoE (#35554)
[docs] Add DFlash2 speculative cells to the Qwen3.8-27B cookbook (#35663)
[Fix] Land the decode mamba checkpoint depth on the tree page under DCP (#35412)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-authored-by: Ke Bao <ispobaoke@gmail.com>
📝 [NPU] Clean up quantization comments (#34829)

Co-authored-by: ronnie_zheng <zl19940307@163.com>
feat(grpc): expose KV event discovery metadata (#35714)

Signed-off-by: Connor Carpenter <connorc@nvidia.com>
TP/PP Consensus checker (#34406)
fix(multimodal): keep LLaVA image fetch off the CPU-preprocess timeout budget (flaky test_mixed_batch) (#35700)
Skip empty linear-attention state buffers in PD transfer (#35689)
[MUSA] Harden CI dependencies and diffusion warmup (#35610)
[diffusion] feat: support out-of-tree models and pipelines (#35713)
[diffusion] feat: let every layerwise component be configurable (#35688)
[diffusion] Refresh eager optimization skills and benchmark safeguards (#35679)
test: switch the Inkling-Small NVFP4 deterministic suite to DSPARK (#35293)
[NPU] [FIX] Fix non-contiguous parameter issue in FIA operator (#34936)
[NPU]Ensure tensors allocated by empty_like are contiguous (#34935)
[Fix] Keep deterministic GDN prefill on Triton (#35632)
[diffusion] quant: support pruned safetensors checkpoints for minimax-h3 (#35418)
[diffusion] feat: plan pinned host memory against the cgroup cap not the machine (#35641)
[Quant] Load compressed-tensors kv_cache_scheme scales (#35455)
[diffusion] feat: add weight source reader (#35668)
[diffusion] CI: add minimax-h3 ref2va audio consistency coverage and guard peak vram (#35511)
[diffusion] feat: support unverified short edge instead of rejecting it for minimax-h3 (#35664)
[AMD][CI] Run Both ROCm 7.2.4 and ROCm 7.2.0 Images on Nightly Test AMD (#35603)

Co-authored-by: Cursor <cursoragent@cursor.com>
[AMD][CI] Default the ROCm 7.2 PR gate to ROCm 7.2.4 Image (#35602)
[diffusion] quant: support gguf (#35370)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Mick <mickjagger19@icloud.com>
Split TRTLLM MHA decode batches by KV sequence length (#34888)
[diffusion] fix: keep large vocab tables in host memory under layerwise offload (#35626)
Fix Grok-2 nightly: derive image-understanding capability from is_multimodal (#33730)
Update deepep for SBO feature (#35450)
[Fix]: exclude SM120 from attn-res TMA dispatch (#35361)

Co-authored-by: 1BIN4 <1741738350@qq.com>
Co-authored-by: L-Ark <fliangae@connect.ust.hk>
Co-authored-by: Chikati <jxudn@connect.ust.hk>
Co-authored-by: mengzili <zilim@ust.hk>
Remove unused MOONCAKE_COMPILE_ARG argument from Dockerfile (#35649)
[HiCache] Allow a retraction host pool smaller than the device pool (#35543)

Co-authored-by: cctry <cctry@fb.com>
Amd/dsv4 shared experts fusion top6 (#32340)

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: kk <43161300+kkHuang-amd@users.noreply.github.com>
Co-authored-by: Thomas Wang <thomawan@amd.com>
[AMD] Add GLM-5.2 MI35x nightly accuracy and perf benchmark (#32570)
update codeowner (#34802)

Co-authored-by: liusy58 <liusy58@smail.nju.edu.cn>
[Docs] Update contribution guide (#35419)
[CI] Surface AMD ROCm 7.2 state in the PR CI-states block (#34813)

Co-authored-by: Michael <michaelzhang-ai@users.noreply.github.com>
Co-authored-by: Chen <bingxche@amd.com>
[Fix] Support 128-aligned hidden sizes in the W4AFP8 DeepEP low-latency requant kernel (#35593)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
[diffusion] UX: report where a component's weights are (#35618)
[XPU] Fix/kimi linear xpu (#34546)

Co-authored-by: github-actions[bot] <github-actions[bot]@users.noreply.github.com>
Co-authored-by: Singh <rohitsi2@iil-login.iind.intel.com>
Co-authored-by: Singh <rohitsi2@iil-gnrap02.iind.intel.com>
[diffusion] fix: keep cosmos3 T=1 fusion on blackwell only (#35612)
[diffusion] CI: use canonical residency selector in nightly (#35615)
[diffusion] UX: reduce per-request log noise (#35614)
[Feature] Add process-local in-memory KV indexer and Router integration (#33370)

Co-authored-by: Wu, Yutong <yutong.wu@amd.com>
Co-authored-by: TianDi101 <ditian12@amd.com>
Co-authored-by: Zhangheng <hzh0425@apache.org>
[XPU][CI] key persistent JIT kernel cache by image content ID (#35337)
[DeepSeek-V4] Add Q8KV8 sparse MLA prefill runtime backend (#32327)

Co-authored-by: Ho-Ren (Jack) Chuang <horenchuang@bytedance.com>
Co-authored-by: Xiaoyu Zhang <1182563586@qq.com>
[misc] Add a comment style rule to .claude/rules (#35597)
docker: fix CUDA-13 build — rename NCCL_VERSION ARG to avoid base image ENV collision (#35587)
[CI][AMD] Run the profiling suite without CUDA graphs on ROCm (#34452)

Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: quitenode <quitenode@users.noreply.github.com>
[AMD] [Docker] Upgrade Python 3.12 + torch 2.11 + triton 3.7 in ROCm 7.2.4 (#30984)

Co-authored-by: Chen <bingxche@amd.com>
[diffusion] fix: reject unsupported modelopt checkpoint algorithms (#35182)
[Spec] Support quantized target lm_head in the DFlash2 selector (#35496)

Co-authored-by: LING ZHI <1747985437lz@gmail.com>
[diffusion] fix: stop reserving nccl device buffers for single-rank groups (#35538)
[AMD] Keep the PTX-inline-asm diffusion norm fusions off on ROCm (fix FLUX warmup crash) (#34481)

Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Make PR babysitter launcher fork-safe (#35575)
Support custom draft worker classes in DSpark (#35397)

Co-authored-by: Yichao Fu <yichaofu@meta.com>
[sampling] Fix int32 offset overflow in top-k renorm Triton kernels (#35571)

Co-authored-by: Xiaozhu Meng <mxz297@gmail.com>
[Kernel] Support wider rows in mega_moe_pre_dispatch (#35372)
chore: bump tilelang to 0.1.12 (#30874)

Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
[Quant] Load compressed-tensors quantized lm_head instead of value-casting it (#35228)
fix(constrained): reject NUL bytes in grammar specs to stop an xgrammar segfault (#34679)

Signed-off-by: Junhao Shen <junshen@nvidia.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-authored-by: Baizhou Zhang <sobereddiezhang@gmail.com>
Co-authored-by: kpham-sgl <khoa.pham@radixark.ai>
[Bugfix] Fix min-new-token EOS handling (#31378)

Signed-off-by: Alexandre Milesi <milesial@users.noreply.github.com>
[HiCache] Simple style change for buffer mode (#35574)
Add docs for TP LMHead optimizaiton (#35283)
Revert "[Feature] Add DeepEPv2 (ElasticBuffer) MoE A2A backend" (#35568)
[Fix] Fix Nemotron-H Mamba illegal memory access under DP attention with CUDA graph (#34561)

Co-authored-by: Brayden Zhong <b8zhong@uwaterloo.ca>
fix(disagg): allow fake transfer with decode DCP (#35409)

Signed-off-by: Alexandre Milesi <milesial@users.noreply.github.com>
feat(openai): Accept the input_audio content part in chat completions (#33606)
[DSA] Trim top-k v2 output modes and tighten its PDL waits (#35041)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
[HiCache] Split the host-memory budget across co-located ranks (#35540)
fix(gemma4): quantize MTP bridge projections (#32440)
[Scheduler] Add configurable decode interval after prefill (#35017)
[Feature] Add DeepEPv2 (ElasticBuffer) MoE A2A backend (#29525)

Co-authored-by: menyu <menyu@nvidia.com>
[Qwen3.5][MTP] Preserve online NVFP4 draft quantization for mixed checkpoints (#35545)
Support Intern-S2-Mobius FP8 (#34908)
[Fix] Support Kimi-K3 ModelOpt mixed NVFP4/FP8 checkpoint (#35077)
[UnifiedTree] feat: support runtime attach/detach (#35269)

Co-authored-by: hzh0425 <hzh0425@apache.org>
[NIXL] Query EP top-k index dtype (#35294)
[Docs] PaddleOCR-VL: update which stage of the pipeline this serves and show real output (#35458)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[diffusion] fix: route quantized vae component repos safely (#35184)
[diffusion] fix: fix multi-group layerwise offload startup memory (#35509)
[Diffusion]  Use current_platform instead of hardcoded "cuda" in cosmos3 guardrails  (#34612)

Co-authored-by: ronnie_zheng <zl19940307@163.com>
[AMD] cookbook: serve Qwen3.5 MXFP4 on MI355X with an fp8_e4m3 KV cache (#35445)
fix: fix transcription & audio-understanding for ASR/audio/speech models (#32611)

Co-authored-by: Singh <rohitsi2@iil-login.iind.intel.com>
[AMD] DeepSeek-V4: route decode wo_a bf16 batched matmul to aiter batched_gemm_bf16 (#33313)

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: Thomas Wang <thomawan@amd.com>
[AMD] DeepSeek-V4 MI355X: eliminate bpreshuffle fp8-scale relayout copy in dense w8a8 linear (#33165)
Add three new test cases (#35502)

Co-authored-by: HeYao <heyao@example.com>
[PD] Deferred decode-side KV release for the NIXL backend (#35360)
[PD] Overlap prefill DP-rank bootstrap queries (#35071)
[HiCache] Support DCP with DSpark (#35221)

Co-authored-by: Cursor <cursoragent@cursor.com>
[docs] Add a fused-kernels page for SGLang Diffusion (#35436)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Qwen3.8-27B Model Support (#34859)

Co-authored-by: Jimmy Shong <69131491+Jiminator@users.noreply.github.com>
Co-authored-by: Brayden Zhong <brayden.zhong@radixark.ai>
Co-authored-by: BBuf <1182563586@qq.com>
Co-authored-by: Zijie Xia <zijie.xia@radixark.ai>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Qiaolin Yu <liin1211@outlook.com>
[AMD][DI][CI] Run MI355X disagg nightly at 7AM UTC (#35467)

Co-authored-by: bingxche <bingxche@users.noreply.github.com>
[diffusion] refactor: gate native encoder quantized checkpoints (#35183)

Co-authored-by: Yiqi Yang <yangyiqi8787@gmail.com>
[CI] Trim the base-c 4-gpu-h100 stage from 5 shards to 4 (#35407)
[Fix] Scale the req_to_token row headroom by attn_dcp_size (#35424)
[AMD] Let the diffusion AITer backend take grouped-query K/V (fix Cosmos3-Nano startup) (#34485)

Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: Michael <michaelzhang-ai@users.noreply.github.com>
VLM: feed the packed qkv projection output to vision backends uncopied (#35336)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[diffusion] chore: reuse shared checkpoint quant metadata resolver (#35174)
[diffusion] optimization: reduce minimax h3 mps memory pressure (#33880)
[Constrained] Support MistralCommon tokenizers in the XGrammar backend (#35215)

Co-authored-by: Xinyuan Tong <xinyuantong.cs@gmail.com>
Fix HiCache PP sync test fixture (#35446)
[Fix] DCP: advertise the logical KV-event block size (#35298)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
feat(api): add sglext_spec (#33518)

Signed-off-by: Muqi Li <muqi1029@gmail.com>
Co-authored-by: Codex <noreply@openai.com>
[HiCache] Batch PP write and load completion sync (#33473)
[Perf] Restore the 16-token router GEMM threshold on SM10X (#34953)

Co-authored-by: Mohammad Angkad <mohammad.angkad@radixark.ai>
[diffusion][Minimax H3]support subblock sparse attention on SM90 (#34680)
[diffusion] fix: make MiniMax-H3 AdaLN cache rebuild transactional (#34993)

Co-authored-by: Mick <mickjagger19@icloud.com>
[HiCache] Buffer-only mode for HiCache host memory layer (#34798)
[diffusion] refactor: reuse srt qwen vision and text modules (#35006)
Fix DP attention on CPU (#12961)
Add fmha_v2 attention backend for SM90/120 (#23112)

Co-authored-by: Yangmin Li <yangminl@nvidia.com>
[diffusion] chore: make --vae-tiling honest, fix the decode oom advice, gate nvfp4 on blackwell (#35353)
quant: extract shared checkpoint quant metadata resolver (#35172)
[diffusion] feat: support cache-dit, cfg gating, attention backend override as per-request param (#35339)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
[perf] overlap page preprocessing, pack the vit, enable prefill CUDA graph for paddle-ocr (#35318)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[Doc] Fix TP and attention-TP group layout in initialize_model_parallel docstring (#34862)

Co-authored-by: NanoByte0513 <167996578+NanoByte0513@users.noreply.github.com>
Co-authored-by: Xiaoyu Zhang <1182563586@qq.com>
[Spec] DFlash2: local convolution + candidate selector (#35371)

Co-authored-by: Jian Chen <jianchen0311@gmail.com>
Co-authored-by: Liangsheng Yin <hnyls2002@gmail.com>
Co-authored-by: hnyls2002 <lsyincs@gmail.com>
install sglang in virtual env instead of system path (#30612)
Apply latest DeepEP branch (#34923)
[Memory] Borrow CUDA graph pool storage for EAGLE sampling (#35375)

Co-authored-by: cctry <cctry@fb.com>
[Fix] Assert the page-aligned SWA evict floor on both PD decode prealloc paths (#35396)
[CI] Skip fast-fail for scheduled stages (#35392)
[Refactor] Share the page-aligned decode alloc lens between EAGLE and DFLASH (#35382)
Laguna: config-driven MoE router scoring (#35362)
[Spec] Page-align the DFLASH decode KV reservation (#35265)
[Fix] Assert the page-aligned SWA evict floor at PD decode prealloc (#35286)
[Fix] Skip padded state slots in the chunked GDN kernel (#33431)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Stop losing Kimi-K3 tool calls to reasoning, constraint conflicts, and truncation (#34881)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-authored-by: Xinyuan Tong <xinyuantong.cs@gmail.com>
Co-authored-by: Xinyuan Tong <115166877+JustinTong0323@users.noreply.github.com>
Fix NIXL cleaner grouping for hybrid cache keys (#35130)

Co-authored-by: Wei Yang <yawei@microsoft.com>
[Metrics] Discount queued prefill load by recent cache hits when waiting-queue matching is off (#35248)
[Fix] Select custom all-reduce v2 by topology capability (#35061)

Co-authored-by: xingyuliu <xingyuliu@fb.com>
Refactor kv cache event mixin into a recorder (#35164)
fix: preserve output logprobs without input logprobs (#34627)

Signed-off-by: jain-ria <riajain@NVIDIA.com>
[diffusion] fix: decouple encoder parallelism from the dit parallel layout (#34713)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
[NPU] Add mxfp4-w4a4 MOE Quantization Support for NPU (#30319)
[PD] Deferred decode-side KV release for aborts mid-transfer (#35049)

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Exclude multimodal-gen NPU jobs from fast-fail cascade (#35238)
[NPU CI] Reorganize test output/log directory structure with workflow context (#33685)
[diffusion] optimization: INT8 Linear + pluggable DiT attention backends for MiniMax-H3 on consumer-level GPUs (#34581)
[Scheduler] Cap prefill-delayer queue target by admission capacity (#35191)
[diffusion] rl: support cosmos3 (#34197)
[kernels] Reorganize ops/diffusion by operator domain behind a lazy facade (#35114)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Add deepseek_v4_flash_w8a8_8p_in32k_out1k_50ms (#35162)

Co-authored-by: HeYao <heyao@example.com>
[AMD] MiniMax-M3 : Fuse QKV+index proj for block-fp8 (#32099)
[AMD] feat(moe): fold padded-topk_ids fill into fused shared-experts append+remap (#31370)
[Rust Server] Add e2e latency metadata and fix Sarashina import (#35125)
test: extend NVFP4 Marlin tests to SM120 (#34327)
Profiling Enhancements [2/3]: detailed execution step annotations (#24911)
[diffusion] chore: filter transformer safetensors by index.json to drop duplicate shard variants (#35107)

Co-authored-by: Emil Bogomolov <zetyquickly@googlemail.com>
[Perf] Hoist DSv4 draft-extend SWA write locs; unify SWA graph buffer naming (#34890)
[Chore] Move version tag helper to release scripts (#35196)
[DSV4] Turn on mhc post pre fusion by default (#35214)
[diffusion] Per-section LoRA adapters on fused linear layers (#34933)
[XPU] Fix decode graph runner is_current_stream_capturing on non-CUDA devices (#35050)
[AMD] Update amd k3 cookbook for PR#34580 (#35263)
[Diffusion][Refactor] Refactor and extract complex RoPE implementation to layers/rotary_embedding for MOVA DiT (#31453)

Co-authored-by: ronnie_zheng <zl19940307@163.com>
refactor: rename chat response token IDs (#35225)
[mem_cache][8/N] refactor: move MambaPoolHost to pool_host.mamba (#31180)
[AMD] Fix Quark Shared Experts Fusion Gate after load-time-override Removal (#35200)
[AMD] Scope the EAGLE greedy-verify TP broadcast to ROCm only (#35195)
[AMD] Add the Kimi-K3 MI35x perf benchmarks in nightly (#34985)

Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: Michael <michaelzhang-ai@users.noreply.github.com>
Skip inkling sheared bias under batch invariance (#35161)
[AMD] Optimize KIMI-K3 with Triton MLA decode kernel by tuning the stage-1 geometry for gfx950 (#34580)

Co-authored-by: Thomas Wang <thomawan@amd.com>
[Docs] Enable PD disaggregation for DSV4 low-latency recipes (#35224)
[DCP] Drop the prefill index-selection syncs by taking each rank's rows by stride (#35084)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
[Feature] Optimize TP LMHead with All-to-All (#32313)
[XPU] Enable fused GDN QKV split Triton kernel on XPU (#30144)

Co-authored-by: Ma Mingfei <mingfei.ma@intel.com>
[CPU] Explicitly import sgl_kernel in CPU kernel tests (#35119)
[diffusion] feat: load quantized H3 text encoder checkpoints (#34986)

Co-authored-by: Yiqi Yang <yangyiqi8787@gmail.com>
[XPU] xpu kernel release workflow (#33679)
[CI] Move DSA PD+MTP+CP Layersplit test to basic B300 test suite (#35220)
docs: sync LMSYS SGLang blog cards (#35218)

Co-authored-by: sglang-bot <sglang-bot@users.noreply.github.com>
Update Qwen3.5 H200 FP8 for AgentX HiCache MTP (#35194)
[Spec] Reduce host-side overhead in ngram draft prep (#35207)
config: one control-plane log for the process (#35028)
config: the readback and the resolving view say what they are (#35027)
config: the per-instance families read the bags (#35026)
config: the DP/EP topology reads come from the parallel bag (#35025)
spec: size the speculative buffers from the bags, not the startup record (#35024)
config: publish before a process reads configuration (#35023)
config: retire the multi-engine accommodation in the runtime context (#35022)
Clean deprecated DeepSeek V4 Environs (#34926)
[metrics] Fix prefill FLOPs estimate to count prefix and per-request causal pairs (#34316)
docs(cookbook): add Qwen3.8-27B DGX Spark configs (#35121)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
[AMD] Add Kimi-K3 8-GPU MI35x nightly accuracy CI (#32568)

Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: Michael <michaelzhang-ai@users.noreply.github.com>
[Misc] Clean up python/sglang package structure (#35062)
[Spec] Support output logprobs with DSpark (#34478)

Co-authored-by: zhisbug <1654062+zhisbug@users.noreply.github.com>
Co-authored-by: Lianmin Zheng <lianminzheng@gmail.com>
[Spec] Relay ngram accept tokens through the FutureMap (#35198)
docs: add NVFP4 quantization option to Kimi-K3 deploy panel (#35168)
[PD] Preserve decode KV across retraction in HiCache (#34801)

Co-authored-by: cctry <cctry@fb.com>
Clean up environ.py: remove dead env vars, unify deprecation handling, move examples to a unit test (#35060)
[diffusion] chore: reuse SRT CLIP encoder blocks (#35004)
Add bit-exact class for MTP (#35143)
[diffusion] chore: reuse SRT SigLIP in Pi0.5 (#34992)
[diffusion] fix: fix h3 swap peft SwiGLU lora_B halves when loading FFN Lora (#34940)
Stabilize GB300 nightly tests (#35044)
[XPU] upgrade sglang xpu backend to PyTorch 2.13 (#31751)

Co-authored-by: MingxuZh <109504044+MingxuZh@users.noreply.github.com>
Co-authored-by: Ma Mingfei <mingfei.ma@intel.com>
[AMD] diffusion: normalize ModelOpt-FP8 weights to e4m3fnuz on gfx942 (#35111)
[AMD] Guard ROCm 7.0 build from using hipMemcpyBatchAsync (#35128)
[AMD] [GLM5] fp8 MLA absorbed bmm for GLM-5.2 on gfx950 (#30519)

Co-authored-by: Thomas Wang <thomawan@amd.com>
Co-authored-by: sogalin_codegen <39478626+sogalin@users.noreply.github.com>
[DSA] Skip indexer KV cache for skip-topk layers (#30531)

Co-authored-by: Brayden Zhong <brayden@radixark.ai>
Co-authored-by: mmangkad <mohammad.angkad@radixark.ai>
[PD] Avoid unused PREBUILT prompt tensor transfer (#35070)
[Spec] Resolve shared-read ends from the backend declaration alone (#35059)
[NPU] Support DeepSeek-V4 DSpark and refactor DSV4 cache management (#33676)

Co-authored-by: JiaruiChang5268 <jc5268@columbia.edu>
Co-authored-by: Kelon <kelonlu@163.com>
Co-authored-by: unknown <z8ruev42yk@gmail.com>
Co-authored-by: Talantan1102 <545811257@qq.com>
Co-authored-by: Talantan1102 <44429302+Talantan1102@users.noreply.github.com>
docs(cookbook): Qwen3.8-27B deployment grid rework (#35065)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Fix sconv track refresh on graph capture (#35042)
[JIT Kernel] Migrate causal_conv1d_fwd and causal_conv1d_update from AOT to JIT (#35031)

Co-authored-by: Mohammad Angkad <mohammad.angkad@radixark.ai>
[Fix] Read the DSA prefill CP flag from the parallel config bag in bootstrap (#35110)
Suppress expected FlashInfer TRT-LLM workspace warnings (#34921)
Fix world-size-one aliasing in MLP batch sync (#34997)

Co-authored-by: wangwenchen0407 <wangwenchen@meta.com>
Fix rope config compatibility and VL/transformers-fallback weight loading (#31575)

Co-authored-by: Claude Sonnet 4.5 (1M context) <noreply@anthropic.com>
Revert "[AMD] [GLM5] Fuse shared-expert append into aiter grouped-topk (skip per-layer append kernel)" (#35105)
fix(hicache): limit load-back pending to write-back (#34519)

Co-authored-by: Zhangheng <hzh0425@apache.org>
[CI] Install sgl-eval in xeon (CPU) Docker image (#34818)
docs: fix Qwen3.8-27B mamba ratio calculator for speculative decoding (#35064)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
[Engine] Freeze GC after server warmup (#34999)

Co-authored-by: Jialin Ouyang <Jialin.Ouyang@gmail.com>
[AMD] Support prefill context parallel two batch overlap for DeepSeek V4 (#33480)
Upd: code owners (#35094)
[DSV4] Emit TMA-aligned UE8M0 scales for FP8 einsum (#34277)
[DCP]Localize HiCache DCP indices once per transfer, not per layer (#34889)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
[Fix] Correct dense FP8 Marlin bias ordering (#35020)
[CPU] Add support for Gemma4 on Xeon (#22498)

Co-authored-by: Copilot <198982749+Copilot@users.noreply.github.com>
Co-authored-by: jianan-gu <jianan.gu@intel.com>
Co-authored-by: Haotong Zou <haotong.zou@intel.com>
[MoE] Add H20 fp8_w8a8 tuned configs for Qwen3.8 (triton 3.7.1) + fix Qwen3_5MoeForCausalLM tuning (#34795)
[Docs] Feature MiniMax-H3 in the popular-models banner (#35068)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
fix tpot by adjusting the sliding max-prefill-size window size (#34856)
[AMD] [GLM5] Fuse shared-expert append into aiter grouped-topk (skip per-layer append kernel) (#31323)

Co-authored-by: Thomas Wang <thomawan@amd.com>
Co-authored-by: HaiShaw <hixiao@gmail.com>
[diffusion] chore: reuse srt siglip vision model (#34988)
[diffusion] Reuse bit-exact modulation fast path for LTX-2.3 (#34930)
[AMD][CI] Add GPT-OSS perf benchmarks to the ROCm 7.2 nightly (#34645)

Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: Michael <michaelzhang-ai@users.noreply.github.com>
[AMD] [GLM5] Skip DSA decode indexer when kv_len <= index_topk (dense k-only fast path) (#31324)

Co-authored-by: Thomas Wang <thomawan@amd.com>
Co-authored-by: HaiShaw <hixiao@gmail.com>
[Spec] Simplify compute_spec_v2_logprobs signature and skip identity gathers (#35058)
[BCG][6/N] Allow prefill breakable CUDA graph for the Kimi archs (#34245)

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Increase post-capture decode memory reserve (#34996)
Add explicit EPLB balancedness reporting modes (#34998)
[Spec] Point multi-layer eagle's last shared-read runner at the draft runner (#35057)
[VLM] Avoid synchronizing multimodal placeholder counts (#34995)

Co-authored-by: Jialin Ouyang <Jialin.Ouyang@gmail.com>
Support unified SWA page mapping in attention metadata (#35000)

Co-authored-by: Yonghao Zhuang <yhzhuang@meta.com>
[Frontend] Apply request header overrides to chat completions (#35001)

Co-authored-by: Ye (Charlotte) Qi <ye.charlotte.qi@gmail.com>
Support model-defined prefill input embedding width (#35002)

Co-authored-by: Lu Fang <30275821+houseroad@users.noreply.github.com>
[Spec] Support logprobs with DSpark speculative decoding (#34696)

Co-authored-by: QAQEthan <QAQEthan@users.noreply.github.com>
Build Rust extensions on demand in source checkouts (#34994)
Clean up playground scripts and add PR babysitter launcher (#35018)
[misc] Rename shared-read boundary to shared-read ends and fix wrapper delegation (#34982)
Add bit-exact guard for extra_buffer_lazy (#35030)
[diffusion] CI: tighten NVIDIA perf baselines (#35016)
[diffusion] Accelerate Cosmos3 T2I QKNorm+RoPE (#34932)
[diffusion][kernel] Accelerate Sana BCG with bit-exact conv post-processing (#34928)
vlm: cache kimi-k3 per-image processor artifacts (#34404)
vlm: streamline vision sdpa reshapes (#34991)
[diffusion] Accelerate lossless Ideogram norm post-processing (#34931)
[diffusion] Enable breakable CUDA graphs for LTX-2.3 (#34929)
Add skill for babysitting PR CI (#35015)
[Quantization] Fix GPTQ scheme attachment broken by LinearBase.scheme default (#34962)

Co-authored-by: Mohammad Angkad <mohammad.angkad@radixark.ai>
[diffusion] chore: refresh docs, retire stale knobs, and fix nightly attribution (#34663)
[diffusion] chore: speed up minimax-h3 vae decode on 2×h100 (#34817)
refactor(hicache): flatten L2 transfer execution (#34793)

GB300 test fails unrelated
[JIT Kernel] Migrate moe_topk_softmax from AOT to JIT (#34509)

Co-authored-by: Xiaoyu Zhang <1182563586@qq.com>
[NPU] Add mxfp4-w4a8 MOE Quantization Support for NPU (#30318)
[AMD] Qwen3.5: guard attn layers against empty DP-attention batch (#34474)

Co-authored-by: jacky.cheng <yichiche@amd.com>
[Fix][AMD] MoRI EP: drop record_stream in TBO dispatch/combine (HSA out-of-resources) (#32746)

Co-authored-by: billishyahao <bill.he@amd.com>
Co-authored-by: Duyi-Wang <duyi.wang@amd.com>
[HiCache] Optimize LogicalHostPool free-list release (#33998)
[AMD][Fix] Qwen3.5: guard zero-grid launch in fused_qk_gemma_rmsnorm(_with_gate) (HIP invalid configuration on idle DP rank) (#31794)
Fix swa eviction frontier for bigram keys (#34870)
[diffusion] refactor: unify component residency controls (#34736)
[AMD][Quantization][Bugfix] Fix bug related to fp8 max on gfx95x for per-token-group quant (ROCm) (#30900)
[AMD] [GLM5] Enable dense-MHA short-context prefill fallback on gfx950 (#30808)

Co-authored-by: Raiden-Makoto <Raiden-Makoto@users.noreply.github.com>
[diffusion] refactor: route minimax h3 vae attention through native backends (#34949)
[diffusion] chore: use native hunyuan3d paint and delight models (#34980)
[diffusion] chore: use native ernie prompt enhancer (#34951)
[diffusion] chore: use native qwen3-vl vision encoder (#34945)
[diffusion] chore: use native qwen2.5-vl generation (#34896)
[Spec] Support mamba-radix-cache-strategy extra_buffer_lazy with DFLASH (#34763)

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
[CI] Pin the allowed_media_domains supplied-instance reads in the step-12 ratchet (#34961)

Co-authored-by: Mohammad Angkad <mohammad.angkad@radixark.ai>
[AMD] perf(sgl-kernel): default block_quota=16 for MLA page_first KV gather… (#30024)

Co-authored-by: Niko Ma <nima@amd.com>
Co-authored-by: figo <fizhang@amd.com>
Co-authored-by: AMD-yanfeiwang <yanfei.wang@amd.com>
[misc] Rename the WAR read-done fastpath to shared-read-done (#34916)
[AMD] Add concat_and_cast_mha_k_pad_kernel to support 12-head and enable K3 aiter prefill kernel (#34837)
Fix dsv4 kl test timeout (#34963)
Add --http2-max-concurrent-streams server arg (#34796)

Co-authored-by: Yilong Zhao <74357408+happierpig@users.noreply.github.com>
update codeowner (#34866)
Fix Whisper transcription for audio over 30 seconds (#33604)
[diffusion] doc: define native diffusion model integration contract (#34952)
[diffusion] model: support ltx-2.5 (#34471)

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Co-authored-by: Mick <mickjagger19@icloud.com>
[diffusion] chore: scope attention backend fallback (#34891)
Fix Python packaging shadowing in DeepEP wheel builds (#34937)
test: restore GLM-4.1V nightly latency threshold (#34811)
feat: add safeguards for remote media URLs (#34892)
[diffusion] Bound overlong weight lock filenames (#34825)
Update sgl-deep-ep release workflow for DeepEP v2 (#34914)
[Kimi-K3] Use explicit SiTU activation for MegaMoE (#34883)
test(step-12): state the bag contract as what resolution produced, and the skill rule that goes with it

`test_bag_values_match_server_args` asserted `bag == field`. That holds today
only because construction resolves in place; step 12 keeps the record raw, and
the plan doc calls this test out as one that becomes **false by design** for
every field resolution fills in.

Rewritten against the resolved projection, which is the half that survives: the
bag carries what resolution produced. The `bag == field` assertion stays as one
line at the end, labelled as the tripwire -- when it starts failing for a
resolution-written leaf, the flip has landed and the bag is the only place the
effective value lives.

The reference is an independent resolution of the same raw input (a fresh,
never-published record) rather than `resolved_server_args_dict()`, which reads
`vars(server_args)` back and therefore only restates the published instance.
And the record goes through the real pipeline on a real mini config, published
through `publish()`: the dummy-model path returns at the dummy boundary with
every sampled leaf still raw, so the old comparison was raw==raw and vacuous
(both Codex catches). Reproducibility (#34094) licenses the sibling as a
stand-in for the pipeline output.

The sample admits only leaves resolution writes on this input on both CI
device shapes (attention_backend, page_size, chunked_prefill_size,
mem_fraction_static), and the raw-differs guard asserts it per leaf -- a
default-count threshold let supplied inputs like `model_path` (no dataclass
default, so any path "differs") stand in for resolution work. Passthrough
leaves (host, hicache_ratio, moe_runner_backend, model_path) move to a
separate projection smoke that claims only what it checks: publish projected
an unchanged field into its namespace. Between the two resolutions the test
restores environ and the EnvField none-flags, so the sibling resolves the
same pristine input rather than the first resolution's leftovers.

And the class runs its body exactly once, like the other dual-resolve
harnesses: a CI retry re-enters after the first attempt leaked process
state, which is the hazard the pristine snapshot exists to rule out.

docs(skill): a supplied-instance read is not automatically safe

The whole-object rule said "keep the supplied-instance contract; don't rewrite
the parameter reads unless the field is runtime-mutated". That is the right rule
for the *object* and the wrong stopping point for the *field*: after step 12 the
record carries the user's raw input, so `server_args.page_size` inside a
runner-owned constructor reads the CLI default rather than the effective value.

The rule now names that second case as step-12 debt with a guard attached
(`test_supplied_instance_exposure_ratchet.py` fails on a new pair, so the
decision is made when the read is written), and names the two shapes that stay
parameter-form on purpose: a helper the resolution pipeline calls with a
`resolved_view`, and a factory whose contract is "build X from the record you
are handed".
config: pin the step-12 debt on the supplied-instance surface

A callee that takes `server_args` keeps the supplied-instance contract, so no
ratchet counts its reads -- and that is right for the *object*. What it does not
cover is what the object will carry after step 12: the instance stays at the
user's raw input, so a callee reading a field resolution fills in starts seeing
the CLI default instead of the effective value.

Measured, not guessed: **297 distinct file/field pairs read one of the 125
fields resolution can write** -- what remains after the earlier members'
conversions (the census found 314 across the package; the page_size,
chunked_prefill_size and graph/limit families were converted by the members
below this one, so the list lands at the remaining debt with no churn). The
census counts three spellings of the read: `server_args.field` off the
parameter, `getattr(server_args, "field", default)` with a literal name, and
the *parked* form -- `self.x = server_args` in a method that takes the
parameter, read as `self.x.field` anywhere in the class. Parking under a
different object, a container, or a computed name stays invisible, like in
every census of this family.

The written-field set is derived in-test from resolved configs against the
dataclass defaults -- the same matrix the context repo's audit tool uses -- and
the union is only as complete as the matrix: fields a matrix entry passes in
are excluded, so each entry must let resolution make the decision the entry is
about. The DWDP shape is the loudest example: `_handle_dwdp` writes `dp_size`,
`enable_dp_attention`, `ep_size` and friends itself, and without a
`{tp_size: 2, dwdp_size: 2}` entry the whole DP/EP topology family (37 pairs
across the launcher, the controllers, the tokenizer and the spec workers)
never entered the written set at all. Multi-item scoring is the same shape in
miniature -- `_handle_multi_item_scoring` writes `disable_radix_cache` itself,
and the `{enable_mis, attention_backend=flashinfer}` entry (the backend is
passed because the handler asserts rather than switches) pins the radix-cache
builder and its friends.

The written set carries **may-write semantics**, statically collected from
every mechanism that can put a resolved value on the record, unioned with the
matrix (construct-and-diff still catches value-level writes statics cannot
name). The enumeration approach kept losing to review -- the DFLASH hook hole,
then the declarative registry (`MODEL_OVERRIDES` forces dtype for two arches
through a setattr applier no assignment scan sees), then the deprecated-alias
loop that writes through a name tuple -- because each round found one more
*mechanism*, not one more field. So all of them are collected now: hook
assignments under `arg_groups/`, the record's own method assignments (the
mooncake layout rewrite, the deepseek-EP mode defaults, the seed fill that
only fires when the caller did NOT supply one -- which construct-and-diff can
never see, since measuring requires supplying), the declarative override
registry, the alias-normalization tuple (drift-guarded), and the
late-resolution keywords. A statically-collected write site that can never
fire is a dead branch to delete upstream, not a reason to shrink the census
(maintainer's ruling). And the pin is split by host: `_EXPOSED` is asserted everywhere,
`_EXPOSED_CUDA_ONLY` (empty today) is where a capability-gated write's
readers go -- one shared exact list cannot hold such a pair at all, since
pinning it fails CPU as "gone" and omitting it fails CUDA as "new".

The resolve-once shape also undercounts on axes a construction call never
sees (each of these was a review catch, and each added pinned families):
resolution branches on the *environment*, so every matrix entry resolves under
the plain env and under the CI shape (`SGLANG_IS_IN_CI`), with the pristine
process state (environ plus the `EnvField` descriptor flags) restored between
entries -- that is where `soft_watchdog_timeout`'s four readers come from. Some
fields resolve *late*, at validation rather than construction
(`declare_late_resolution`): those writers are collected statically by keyword
-- `lora_paths`, `reasoning_parser`, `tool_call_parser` -- with the one dynamic
`**detected` site spelled out in a table guarded against drift. Fields holding
only a `default_factory` are materialized rather than skipped, and
`tokenizer_path` / `served_model_name` -- always filled from `model_path` --
leave the passed-inputs exemption and pin their twelve readers. A module the
census cannot parse fails the test instead of shrinking it, which immediately
caught a BOM-carrying file every previous census had silently skipped (the
scans read `utf-8-sig` now). And the test registers on the CUDA runner besides
the CPU suite, because capability-gated writes only open on real hardware;
AMD is intentionally not registered -- an exact pin cannot be verified from
any pinning host -- with the reasoning in the header.

**A second axis is already wrong today**, independent of step 12. Some config is
decided after publish and recorded with `get_context().override(...)` -- elastic
EP resizing `ep_size`, a weight update rewriting `model_path` / `load_format`,
HiCache attach naming a storage backend, adaptive speculative decoding moving
`speculative_num_steps`. That write reaches the bags and never the record, so a
supplied-instance read of one of those fields answers with the startup value from
the moment the override lands. **73 pairs over 13 fields** are in that position -- including the overrides
that arrive as `**kwargs`: the collector statically resolves dict-literal
expansions (the HiCache attach shape, whose write/read pairs on
`hicache_write_policy` / `hicache_storage_prefetch_policy` were invisible
before) and fails loudly on anything it cannot resolve, with
`update_server_args` exempted by name because its key set is the API's
caller's, not this file's.
Whether each is a defect depends on ordering -- a value copied at construction,
before any override, is fine -- so the axis is pinned as a measurement with the
same growth guard, not as a list of bugs.

One of them *was* a defect and is fixed at the base of this stack: the
linear-attn dispatch table rebuilt itself from the record after the SM100 GDN
prefill decision had been recorded in the bag, so a second runner's rebuild
dropped it. That choice is a per-runner stamp now and is not recorded
process-wide at all, which is why neither the read nor the field appears on this
axis.

The list is pinned both ways, on both axes. A new pair fails, because the moment
to decide where a resolved value comes from is when the read is written, not
during the flip; a disappeared pair fails too, naming the entry to delete, so the
registry stays a measurement rather than a memory of one. Both axes
reverse-verified: a new read of a written field is reported by file and field.

Per-field dispositions live in the plan doc; several of these are "should this
callee take a config at all?", which is a design call rather than a sweep.

test(step-12): tripwire on the EPD guard that a raw record would silence

`_reject_missing_dispatched_encoder_embedding` is one of the two reads the
step-12 audit calls a blocker: it keys on `encoder_transfer_backend`, a field
resolution fills in, off a handed record. Today that record carries the
resolved value; after the flip it stays at the argument default `"auto"`
(`ENCODER_TRANSFER_BACKEND_CHOICES[0]`) for every auto-resolved launch and the
503 stops firing -- a guard that goes quiet, which no existing case notices.

The tripwire resolves a real language-only Kimi-K3 TP2 launch (a mini config,
the shape whose auto pick is `"zmq_to_tokenizer"`) and asserts the guard
rejects with the record resolution produced. A fixed double cannot trip on the
flip -- it would keep handing the guard the resolved value by construction --
so the record has to come from resolution itself: when step 12 lands, this
same launch hands the guard `"auto"`, the rejection silently stops, and this
test fails, which is exactly the signal that this reader needs the resolved
value from somewhere else (the per-engine overlay or the bag).

The launch pins `mamba_radix_cache_strategy=no_buffer` (+ the overlap-off it
requires): resolution's hybrid state-cache sizing branches on the host device
and asserts a GPU stack for extra_buffer, which a CPU CI runner does not have,
while the guard under test reads a field independent of that branch. The case
restores env *and* the EnvField descriptor flags -- a real resolution leaves
state os.environ does not carry.
config: the post-publish consumers of the supplied-instance surface read the bags

config: the speculative workers take page_size from the bags

Seven worker constructors stored `self.page_size = server_args.page_size` off
the handed record. They all run after publish and all keep a copy of a
process-level value, which is the first row of the plan doc's supplied-instance
disposition table -- so they read `get_schedule().page_size`, and a post-publish
override now reaches them like it reaches every other consumer.

The supplied-instance census named the seven pairs; the exposure ratchet in the
next member pins what remains after this batch of conversions.

config: the post-publish chunked_prefill_size consumers read the bags

Four of the ten supplied-instance `chunked_prefill_size` reads are plain
post-publish consumers -- the EPLB recorder's buffer sizing, the deep-gemm
compile warmup (five reads), the KV-cache builder's effective size, and the
ngram embedding manager's assert. All are reached from runner init, so they read
`get_schedule()`.

Two are deliberately left: `create_kt_config_from_server_args` builds a config
*from a supplied record* by name and contract, and `CanaryLaunchCapacities.from_args`
is the same shape. Converting those would change what the function is, not where
it reads -- the plan doc's disposition table says so per field.

config: the remaining post-publish graph/limit consumers read the bags

Three more of the census's supplied-instance debts are plain post-publish reads: the dspark worker's
cuda-graph decode sizes, the dspark planner's SPS table bound
(`max_running_requests`), and the LoRA manager's cuda-graph moe buffers. The
dspark worker is the clearest of them -- it already read
`get_exec().graph.cuda_graph_config.decode.bs` thirty lines below the instance
read, so the file disagreed with itself about where the same value comes from.

Left where the function's contract is "build a config from the record you are
handed" rather than "read this process's config":
`create_kt_config_from_server_args`, `DllmConfig.from_server_args`,
`CanaryLaunchCapacities.from_args`, `build_compilation_config`. Changing those
would change what the function is.

config: the runner, scheduler and offload manager take page_size from the bags

The same `self.page_size = server_args.page_size` shape as the speculative
workers, in the three remaining process-owned constructors: `ModelRunner`,
`Scheduler`, and the decode-side KV offload manager. The scheduler process
publishes before any of them run. The one path that did not is `ModelRunner`
constructed standalone -- `python -m sglang.benchmark.one_batch` and the manual
runner tests build it with no prior publish, and the constructor's own publish
sat below this read -- so that publish moves above the constructor's first bag
read instead of leaving a window where the runner half-exists unpublished.

Left where the read belongs to something else: `utils/common`'s predicates are
called only from the resolution pipeline with a `resolved_view`,
`allocation_sizing` takes the config its callers supply by contract, and
`CudaVmmFeatureTransport` is tokenizer-owned -- one per tokenizer worker, which
is the per-instance boundary.

The conversion left the offload manager parking a record it no longer
reads; the parked copy goes with the read (the constructor parameter stays
-- its hicache sizing still reads it directly).
config: the alias form of the runner-side instance read

The previous batch counted `self.server_args.X` and called the runner surface
done. It was not: the same read spelled through a local alias --
`server_args = model_runner.server_args` (or `sa = kvc.server_args`, `args = ...`)
followed by `server_args.leaf` -- is the same process-global read wearing a
local name, and the AST census counts **57 of them** across eleven files that
the grep never saw. Census per function, following the alias.

52 were leaves and go to their bag (`spec` 11, `schedule` 9, `memory` 7,
`exec.graph` 5, `exec.moe` 5, `parallel` 4, `disagg` 4, `model` 3,
`exec.mamba` 2, `exec.overlap` 2). Five were not leaves:
three derived members on the eager runner --
`max_speculative_num_draft_tokens` and `enable_mamba_extra_buffer` already had
accessors, and `max_prefill_buffer_tokens` gets one (all its inputs are `schedule`
leaves plus the configured PP size, so it derives from the bags and follows a
post-publish override; `TestDerivedPredicatesAgreeAcrossTiers` pins it against
the member over a 48-case matrix) -- plus `get_attention_backends()`, which the
same commit routes through `attention_backends()`, and a dict that merely shares
the name (`server_args_dict.items`). That dict is the one read left behind.

`build_attention_backends` also stops resolving the pair from the record: it
runs after publish, so it asks `attention_backends()` like every other consumer.
The draft override on the runner still wins first.

`dispatch_event_loop`'s three PP checks read the *configured* PP size, not the
live topology: the MLX runner stub never initializes torch.distributed, so the
live property asserts before the MLX event loop can start (a Codex catch). The
configured leaf answers the same value wherever the live groups exist.

`flashinfer_gdn_prefill_default`'s guard is the one read here that asks what the
*operator* named rather than what the config resolved to, and the bag leaf now
answers exactly that: the per-runner auto-default is stamped on the runner and
deliberately never recorded process-wide, so nothing writes that leaf after
launch and reading it back cannot mistake another runner's default for a flag.

Three test doubles injected a `SimpleNamespace`/`MagicMock` record for exactly
these reads and now publish instead (pool configurator, cache registry, GDN
prefill policy) -- the fixture publishes what the case configures and hands the
published instance to the whole-object contracts that still take one.

The functions this sweep partially converted stop mixing sources (review
catches): the flash-attention constructor's remaining seed reads
(`speculative_eagle_topk`, `speculative_algorithm`, both deterministic gates)
read their bags next to the leaves already converted;
`_should_disable_scheduler_metadata_precompute` reads the parallel config
leaves itself instead of taking the record (its alias binding was the last
use); and the autotune gates (`disable_flashinfer_autotune`, deterministic,
`flashinfer_autotune_skip_ops`) join the moe leaves the same function already
reads from the bags. The pool-configurator fixture drops a parameter nothing
published or read.
config: spell out the one dynamic config read the census could not see

`_is_dsa_active` asked `getattr(server_args, "_is_dsa_model_arch", False)`, and
that name has never existed on `ServerArgs` -- it arrived as a placeholder with
the CP strategy abstractions (#27313), so the getattr default has always decided
the predicate. A dynamic read of a name nothing sets is the one shape the config
census cannot follow, and it looked like a live decision while being dead.

Spelled as the constant it evaluates to, with the placeholder written down: what
it should ask (whether this process runs a DSA model arch) is the CP path's
call, and its only consumer, `ContextParallelStrategy.per_layer_attn_cp_comm`,
has no readers yet.

That was the sole entry in the read ratchet's `_INERT_DYNAMIC_READS`, so the
exemption list is gone with it -- there is no way to exempt a read from the
baselines any more, which is the invariant worth having. The `counted()`
indirection it existed for goes too (verified the three shapes it guarded still
report: direct, `getattr`, and an attribute-parked alias).
config: decisions keyed on the attention backend read the configured pair

`--attention-backend` is one field of three: a launch that sets only
`--prefill-attention-backend` or `--decode-attention-backend` leaves the base
field at `None`. Seven decisions read that base field alone and therefore
answered from a field the operator never set. `attention_backends()` is the
pair with the base-field fallback already applied, so each site now asks it for
the half it actually needs:

- `inkling_common/attn` assembles backend-specific kwargs (rel_bias / score
  mods) and gates its fused prologue; the backend those describe is the one
  `self.attn` dispatches to, so `serving_attention_backend()` selects the pair
  member by `forward_batch.forward_mode`, mirroring
  `HybridAttnBackend._select_backend` exactly -- draft-extend routes through
  the prefill branch like the dispatcher does -- and preferring the
  runner-stamped pair, so a draft runner answers with its own backend. That
  preference only works if every backend that can enter a ForwardContext
  carries the stamp, so `DraftBackendFactory._create_backend` now stamps its
  products with the backend it resolved (draft override first), and the
  draft-extend conv-sidecar wrapper copies the wrapped backend's stamp -- the
  replacement backends the spec workers install had no stamp at all and fell
  back to the target's configured pair.
- The chunked-prefix-cache gate is a *prefill* feature -> prefill half. Reading
  the base field switched the feature off for every prefill-only configuration.
- `init_deterministic_inference_config` maps *prefill* knobs
  (SPLIT_TILE / PREFILL_TRUNCATION_ALIGN) -> prefill half; the map missed and
  left truncation unset.
- `two_batch_overlap` computes extend positions -> prefill half.
- mrope's interleaved-rope kernel runs in both phases -> both halves must
  support triton. This one is not conservative when it misreads:
  `support_triton(None)` answers **True**, so a `--prefill-attention-backend
  torch_native` launch took the triton path.
- The req-to-token writer has one caller, `alloc_for_extend` -> prefill half;
  its fallback pays several `.item()` syncs per request, so gating it on the
  decode half too would send every extend of a mixed launch through the slow
  path. `get_last_loc` (the spec-decode allocator's helper) keeps the
  both-halves reading: verify tokens are served by either half depending on
  `speculative_attention_mode`.
- The flashinfer version floor is a guard; it never fired for a launch that
  pinned flashinfer through a split field.

One more site the census found is not converted here: `gpt_oss` derives its
`sinks` parameter dtype from the backend, and a single parameter dtype cannot
serve a split pair (FA4 asserts bfloat16, trtllm_mha consumes float32), so
that one is a behaviour question rather than a config-source one and is fixed
in its own PR.

`test_split_attention_backend_decisions.py` pins the callable decisions by
calling them under a split-only publish, and pins the remaining ones
statically -- the file/why map fails if any of them goes back to the base field
(reverse-verified). It also asserts the `support_triton(None) is True` trap the
sweep exists for.

The stamp comes from the constructor, not the request: every factory leaf
answe…
hanwlax pushed a commit to hanwlax/sglang that referenced this pull request Aug 28, 2026
… k-only fast path) (sgl-project#31324)

Co-authored-by: Thomas Wang <thomawan@amd.com>
Co-authored-by: HaiShaw <hixiao@gmail.com>
jakki-amd pushed a commit to jakki-amd/sglang that referenced this pull request Sep 9, 2026
… k-only fast path) (sgl-project#31324)

Co-authored-by: Thomas Wang <thomawan@amd.com>
Co-authored-by: HaiShaw <hixiao@gmail.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants