Skip to content

[Fix] [NVIDIA] modelopt: Keep SM12x FP8 batch-invariant in deterministic mode - #44861

Open
yibol9768-alt wants to merge 3 commits into
sgl-project:mainfrom
yibol9768-alt:fix/modelopt-fp8-batch-invariance
Open

yibol9768-alt wants to merge 3 commits into
sgl-project:mainfrom
yibol9768-alt:fix/modelopt-fp8-batch-invariance

Conversation

@yibol9768-alt

@yibol9768-alt yibol9768-alt commented Oct 17, 2026 •

Copy link
Copy Markdown

Motivation

On SM12x, ModelOpt static-FP8 linears can change reduction schedules between cold prefill and a shorter prefix-cache-hit batch despite deterministic inference being enabled. The same input sequence then receives different scores.

Related to #41351's ModelOpt path. Credit to furedericca-lab for the diagnosis and initial fallback, and twu3202 for identifying the small-M gap. This does not claim to resolve every hybrid/GDN issue in that thread.

Modifications

  • In SM12x deterministic mode only, reuse Triton scaled MM with fixed N=64/K=128 reduction tiles and M-only tiles 16/32/128. Normal-mode and forced-Marlin dispatch remain unchanged.
  • Zero masked A/B loads for partial K tiles.
  • Extend the existing registered SM120 test file with 12 cached-suffix comparisons and two CUDA Graph replay cases covering bias, strided 3D inputs, BF16/FP16 and a K=80 tail.
  • The 2026-10-22 refinement changes only the deterministic GEMM's pipeline from two stages to three, retaining the same tiles, warps 4, quantization, arithmetic and scaling. It overlaps loads while preserving the reduction schedule.

Accuracy Tests

Validated 2026-10-22 on one RTX 5090, 400W, TP=1, using the complete official SGLang 0.5.23.dev20261020+g55b2aef7ad nightly: PyTorch 2.14.1, Triton 3.8.0, FlashInfer 0.7.2, sglang-kernel0.4.10 and its official Rust radix cache. Complete source files are installed before each process starts; normal package imports are used, with no function/AST injection. This is official-nightly validation, not a latest-main rebuild or upstream CI.

Model: nvidia/Qwen3-8B-FP8@2cebc4c89e25abc17668c81b01dceaf3d8b914d5, BF16 activations, automatic FP8 KV, deterministic mode and Triton attention.

Check Result
Complete existing SM120 test file 137 passed
Additional two-stage/three-stage tensor comparisons, BF16/FP16, two seeds, eight N/K shapes, M boundaries, bias, strided input and tails 288/288 bit-exact, also exact against full-batch suffixes
Real model: prefixes 128/1024/4096, batches 1/4/2, two repetitions; two Graph boots and one eager boot per PR arm 18 paired cache cases exact; cold/cached scores and warm repeats have zero delta; selected logprobs also match across pipeline versions
Complete fixed-work outputs, including warmups and measured repetitions 180/180 texts and token sequences exact between pipeline versions
Complete GSM8K test rows 32–159, cap 512, Graph mode 128/128 texts and token sequences exact between pipeline versions;120/128 correct in each, no missing final answers

The 128 questions reuse previous data and test preservation of the existing PR's outputs. They are not an independent quality evaluation, an accuracy gain, or evidence of statistical equivalence to upstream. Earlier upstream/candidate checks remain disclosed:15/16→16/16 on a16-question nightly subset, and119/128→117/128 on the older mixed runtime. See the prior nightly results and historical128-question results.

The unmodified nightly FP8 path still shows nonzero cold/cache selected-logprob deltas in the fresh Graph control. An independent RTX PRO 6000 Max-Q / Qwen3-14B-FP8 check also corroborates the fix at prior head 5f90880cc2; it predates this stage-only refinement. The original hybrid TP=2 model, SM121 and torch.compile have not been tested locally.

Raw paired results, full generated outputs, tests, source snapshots, scripts and hashes.

Speed Tests and Profiling

The isolated probe identifies GEMM as the main component cost in these cases; increasing pipeline depth is faster for several projection shapes while all 64 explored cases remain exact. These cache-evicted microbenchmarks are diagnostic; the end-to-end results below determine the performance claim.

32 forced output tokens per request, one warmup and five measured samples per cell per boot, cache flushed before requests, passive cooling to<=55C outside timing. PR-arm Graph boot order is A(two stages), B(three stages), B, A. Table medians pool ten samples per PR arm; per-boot medians, ranges and every sample are retained. The upstream Graph reference and each eager arm have one boot/five samples. Graph capture covers decode 1/2/4 and prefill 4/128/256/1024/2048/4096/8192; unsupported partial prefills may fall back to eager.

Prefix / batch Graph two stages → three Wall-time reduction Upstream Graph reference Eager two stages → three
128 / 1 0.3119s → 0.2796s 10.3% 0.2684s 1.4868s → 1.4839s
128 / 4 0.3539s → 0.3251s 8.1% 0.3322s 1.5320s → 1.5275s
2048 / 1 0.5314s → 0.5143s 3.2% 0.4766s 1.5629s → 1.5642s
2048 / 4 0.8900s → 0.8688s 2.4% 0.7953s 2.0843s → 2.0747s

The stage-only update reduces this Graph matrix's pooled median wall times by 2.4–10.3%, most consistently on the short prompts. The 2048/1 cell varies between boots: current medians0.5285/0.5895s, candidate0.4982/0.5434s. Eager differences are within about 0.5%, so no meaningful eager speedup is claimed.

Against the single fresh upstream Graph reference, the three-stage path ranges from 2.1% lower to9.2% higher wall time. This does not eliminate the determinism/cost tradeoff or establish a general speedup. Earlier 5.7–25.8% nightly and 5–20% mixed-runtime Graph costs describe the prior two-stage measurements and remain in their historical evidence. All ten successful jobs stayed below 72 C, peaking at 66C with the 400W limit unchanged.

Full samples, boot medians and measurement limits.

Duplicate check: inspected #34429 (small-M dispatch), #42865 (block/MXFP8), #38921 (SM89 prequantized inputs) and #44076 (GDN chunk alignment); these do not implement this static per-tensor deterministic path. Original pinned diff audit.

AI assistance was used for diagnosis, implementation, validation and this description. The human contributor reviewed the original fix and authorized this performance investigation and PR update.

Checklist

  • Applicable pre-commit hooks pass.
  • Regressions extend the existing registered file.
  • Real pretrained-model correctness, complete outputs and performance checks are recorded with limitations.
  • No new public interface; deterministic-mode behavior and tradeoffs are documented here.

Review and Merge Process

Maintainer review and upstream CI follow the normal project process. Local validation and independent reproductions do not constitute upstream CI or maintainer approval.


CI States

Latest PR Test (Base): ❌ Run #39746823806
Latest PR Test (Extra): ❌ Run #39746823720
Latest PR Test (AMD ROCm 10): ❌ Run #39746823779

…tic mode

Co-authored-by: OpenAI Codex
Signed-off-by: Yibo Liu <240359763+yibol9768-alt@users.noreply.github.com>
@yibol9768-alt
yibol9768-alt force-pushed the fix/modelopt-fp8-batch-invariance branch from ff5adcd to f1cae93 Compare October 17, 2026 06:51
Keep N/K reduction tiles fixed while using row tiles suited to decode and
prefill. Cover biased strided BF16/FP16 inputs under CUDA Graph replay.

Co-authored-by: OpenAI Codex
Signed-off-by: Yibo Liu <240359763+yibol9768-alt@users.noreply.github.com>
detain pushed a commit to detain/sglang that referenced this pull request Oct 20, 2026
…roject#44861, part 5a)

Port the deterministic-gated half of PR sgl-project#44861 (net diff of the two upstream
commits f1cae93 + 5f90880 from the fork-point 9c7b338, excluding
kernels/ops/gemm/fp8_kernel.py, which lands separately as part 5b).

Why this stack wants it: our ModelOptFp8LinearMethod.apply() carried ZERO
batch-invariance code, so with a ModelOpt FP8 (or mixed NVFP4) checkpoint on
SM12x the same rows score differently in prefill vs decode.  That is the exact
class of divergence the prefill-vs-decode KL consistency probe measures, so the
NVFP4 arm could never be checked with --enable-deterministic-inference.

Landing (3 files, +120/-1, `git apply -3` from the squashed net diff -- all
three hunks applied CLEANLY, zero conflicts):

  * fp8_utils.py: new apply_fp8_linear_deterministic() (+33), placed exactly at
    the PR's anchor (between use_aiter_bpreshuffle_gemm() and
    apply_fp8_linear_bmm_flashinfer()).  It quantizes with the static per-tensor
    scale and calls triton_scaled_mm() with BLOCK_SIZE_N=64 / BLOCK_SIZE_K=128
    PINNED and only BLOCK_SIZE_M tiered on the row count (16 / 32 / 128), so the
    fp32 accumulation order of every output element is independent of batch
    size: rows are bit-identical between a full prefill and a shorter suffix
    batch (the PR's own test asserts rtol=atol=0).  Our in-tree
    triton_scaled_mm already honours caller-supplied tile sizes when
    use_heuristic=False (the same entry point our tuned-channelwise FP8 path
    uses at fp8_utils.py:2433/:2438), and static_quant_fp8(..., repeat_scale=
    False) keeps x_scale per-tensor, so nothing in the helper's argument
    contract depends on M.
  * modelopt_quant.py: import get_exec alongside get_platform (:68) and
    `apply_fp8_linear_deterministic`, and insert the
    `use_sm120_fp8 and get_exec().deterministic.enable_deterministic_inference`
    branch in apply() between the Marlin branch and the SM12x try_sm120_fp8_linear
    facade -- i.e. deterministic mode deliberately gives up the small-M fast path
    in exchange for a fixed reduction schedule.  Anchors verified in our tree:
    self.use_sm120_fp8 is set at :559 (cuda_capability[0] == 12) and
    get_exec().deterministic.enable_deterministic_inference is the established
    read path here (20+ in-tree users, e.g. fp8.py:1267 and
    moe_runner/triton_utils/fused_moe_triton_config.py:72).
  * test/registered/kernels/ops/gemm/test_sm120_fp8_linear.py: the PR's 14 new
    tests (+74), byte-identical (zero drift on that file).  NO extra gating
    needed from us: the file already ends its import block with
    `if not (torch.cuda.is_available() and get_device_capability() == (12, 0)):
    pytest.skip(..., allow_module_level=True)`, so on this CPU-only box it
    collects as `1 skipped` both before and after (verified).  This box IS
    SM120, so those 14 tests will actually run in a GPU window -> recorded as a
    GPU debt (cache-suffix bit-equality + CUDA-graph replay with bias and a
    strided input).

Fork survival: fp8_utils.py (+41/-163 drift) and modelopt_quant.py (+210/-61
drift) both took the PR hunks at their own anchors; the SM12x facade branch,
sm120_fp8_alpha precompute, _can_use_flashinfer_bmm/use_flashinfer_bmm path and
the entry-190 Marlin/pure-FP8 gates are untouched (hunks are additive).
Behaviour with --enable-deterministic-inference OFF (prod fp8 arm) is unchanged
by this commit: the new branch is gated on the flag.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
(cherry picked from commit 5f90880)
detain pushed a commit to detain/sglang that referenced this pull request Oct 20, 2026
…art 5b)

The second half of PR sgl-project#44861, landed as its own commit because it is NOT gated
on --enable-deterministic-inference and therefore touches a shared kernel.

kernels/ops/gemm/fp8_kernel.py::scaled_mm_kernel is launched from exactly one
place, triton_scaled_mm() (:1112), which has three in-tree callers, all in
srt/layers/quantization/fp8_utils.py: the new apply_fp8_linear_deterministic()
(2275) and the two tuned-channelwise w8a8 paths (2466 / 2471).

Why the explicit `other=0.0` matters: `tl.load(ptr, mask=...)` without `other`
leaves masked lanes "don't care" per Triton semantics, and those lanes are fed
straight into `tl.dot(a, b, accumulator)`.  A masked-out ROW or COLUMN only
pollutes outputs that are discarded anyway, but the K-tail (masks_k) pollutes
VALID outputs: every masked k lane contributes a_garbage * b_garbage to the
fp32 accumulator of every row in the tile, and an uninitialised e4m3 register
can even hold a NaN pattern.  Pinning the masked value to 0.0 makes the
reduction over a ragged K tail exact and reproducible, which is precisely what
the deterministic helper needs (it fixes BLOCK_SIZE_K=128, so any K that is not
a multiple of 128 hits the tail -- e.g. the PR's own 96x80 case), and it also
matches what every other reduction kernel in this file does.

Does any prod behaviour change?  No, on two independent grounds:
  * The masked lanes only exist when K % BLOCK_SIZE_K != 0 (or N % BLOCK_SIZE_N
    != 0, which only affects discarded columns).  The tuned-channelwise configs
    in-tree use BLOCK_SIZE_K in {64,128,256,512,1024} against K in {4096,12288}
    (python/sglang/kernels/ops/gemm/configs/*fp8_w8a8_channelwise.json), all of
    which divide exactly, so those launches mask nothing and compile to the same
    arithmetic.  Prod's fp8 arm additionally pins --fp8-gemm-backend cutlass and
    the only channelwise JSONs present are for NVIDIA_L40S, so
    get_w8a8_channelwise_fp8_config() returns None on this device family today
    and the triton branch is not even taken.
  * Everything else that reaches this kernel is inside the new deterministic
    branch from the previous commit (flag OFF => unchanged).

So for the fp8 prod arm this is a no-op-by-construction change; it becomes
load-bearing the moment a ragged-K shape runs through triton_scaled_mm (the
deterministic NVFP4/FP8 arm, or any future tuned config with a non-dividing K).

Landing: `git apply -3` of the net-diff file selection, CLEAN; zero drift -- the
pre/post blob ids (ff32551..ff7c9f5) are byte-identical to upstream's.
Its tests live in test/registered/kernels/ops/gemm/test_sm120_fp8_linear.py,
which already module-level-skips off SM120 (verified `1 skipped`).

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
(cherry picked from commit f1cae93)
@waynehacking8

waynehacking8 commented Oct 21, 2026 •

Copy link
Copy Markdown

On a second SM12x part, an RTX PRO 6000 Max-Q (SM120) with nvidia/Qwen3-14B-FP8, head 5f90880 makes cold and prefix-cache-hit selected-token logprobs identical in all 12 comparisons, where main differs in all 12 (max abs delta 0.5 to 2.375). The CUDA graph cost matches the 5090 numbers the body gave for that head at batch 1 and is higher at batch 4. All of this is for 5f90880; I haven't run the current three-stage head 6a7695d.

CUDA graph, wall time for 32 output tokens, prefix / batch main, s PR, s PR / main PR / main on the body's 5090 (body for 5f90880)
128 / 1 0.422 0.504 1.193 1.189
128 / 4 0.493 0.577 1.171 1.057
2048 / 1 0.711 0.901 1.268 1.258
2048 / 4 1.167 1.477 1.265 1.169

The comparisons are the body's cases (prefixes 128/1024/4096 at batches 1/4/2, two repetitions) in one eager and one graph boot per arm. Times are medians of five, and the ranges don't overlap in any graph cell. In eager mode PR / main runs from 0.89 to 1.055, but those boots shared the card with another user's job; the graph boots had it alone. The PR's test_sm120_fp8_linear.py passes 137 of 137, and on main's code its cached-suffix cases fail 5 of 12 (4 of 12 on the body's 5090).

Setup: the body's engine settings (deterministic inference, Triton attention, TP1, BF16), fp8 KV, the PR diff applied to main 4fc86f7. Logs and scripts: https://gist.github.com/waynehacking8/e8474d2b9449f657396df5a040566e16

Use three pipeline stages instead of two without changing the N/K tiles,
row tiles, accumulation/scaling order, or SM12x deterministic dispatch.

On one RTX 5090 with the complete official SGLang nightly, the existing
137-case SM120 suite passes and 288 additional tensor comparisons are
bit-exact to the two-stage path. Qwen3-8B-FP8 preserves all 128 complete
answer texts/token sequences and 18 paired cache comparisons.

ABBA graph boots reduce the fixed-work matrix medians by 2.4-10.3%;
eager times are effectively unchanged. Local results are bounded by the
single device/model and do not replace upstream CI.

Co-authored-by: OpenAI Codex
Signed-off-by: Yibo Liu <240359763+yibol9768-alt@users.noreply.github.com>
@adcostanza

Copy link
Copy Markdown

Validated on a hybrid GDN model, which the description leaves out of scope: it fixes this case.

Setup. One RTX PRO 6000 (SM120), TP=1. Image lmsysorg/sglang:nightly-dev-cu13-20261101-cda284e4 with the PR head 6a7695ddc4 merged on top (no conflicts), bind-mounted over the installed package. Model ukisai/Swift-1.5-Qwen3.8-27b-NVFP4 @ 25482027, a ModelOpt mixed checkpoint of a Qwen3.8-27B fine-tune: the GDN in_proj_qkv / in_proj_z / out_proj and the full-attention q/k/v/o_proj are static per-tensor FP8, the MLPs are NVFP4. Flags: --enable-deterministic-inference --attention-backend triton --kv-cache-dtype auto --mamba-ssm-dtype bfloat16 --mamba-radix-cache-strategy extra_buffer_lazy --chunked-prefill-size 2048 --page-size 64 --enable-hierarchical-cache --hicache-size 64 --hicache-write-policy write_through --schedule-policy lpm --enable-session-radix-cache --mem-fraction-static 0.41 --max-running-requests 4.

Result, python3 -m sglang.test.test_deterministic --test-mode radix_cache, cached vs uncached prefill:

cached uncached
nightly cda284e4 -5.0687 -4.8716 FAIL
+ this PR, prefill CUDA graph on -3.7109196186 -3.7109196186 PASS
+ this PR, eager prefill -3.7109196186 -3.7109196186 PASS

A prefill-only probe (cold, prefix hit with a 37-token suffix, cold again; 101 tokens) gives -6.1011090279 all three times with the PR.

Where the difference came from on this model (diagnosed independently before finding this PR, by dumping per-layer tensors for the hit tokens, cold vs hit):

  • The first tensor that differs is layer 0's GDN out_proj output (1.2e-4, from bit-equal inputs). It accumulates in the residual stream and reaches layer 3's norm input as 1.3e-3. NVFP4 MLP outputs stay bit-equal.
  • These FP8 projections go through ModelOptFp8LinearMethod.apply → apply_fp8_linear_bmm_flashinfer → flashinfer bmm_fp8 (use_flashinfer_bmm=True); try_sm120_fp8_linear returns None for these shapes.
  • Standalone, on the real layer-0 out_proj weight (K=6144, N=5120), rows computed in a call of M rows vs the same rows inside a 101-row call: bmm_fp8 differs by 0.0039 at M = 8, 16, 37, 64; torch._scaled_mm and sgl_kernel.fp8_scaled_mm are exact at M = 1, 8, 16, 37, 64, 100.
  • With the PR, all 128 FP8 linear calls per forward (M=101 cold and M=37 hit) take the new use_sm120_fp8 and deterministic branch into apply_fp8_linear_deterministic, which sits before the bmm_fp8 branch, so bmm_fp8 is never reached.

Not measured: throughput of the deterministic Triton path on this model.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

jit-kernel quant LLM Quantization

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants