Skip to content

[AMD] Add fused all-reduce RMSNorm per-token FP8/MXFP4 quant - #29723

Open
mqhc2020 wants to merge 23 commits into
sgl-project:mainfrom
mqhc2020:marv/ar_norm_per_token_quant_fusion
Open

mqhc2020 wants to merge 23 commits into
sgl-project:mainfrom
mqhc2020:marv/ar_norm_per_token_quant_fusion

Conversation

@mqhc2020

@mqhc2020 mqhc2020 commented Jun 30, 2026

Copy link
Copy Markdown
Contributor

Motivation

The existing fused AllReduce → RMSNorm → quant path only supports per-group FP8 scales.
Checkpoints that use per-token FP8 activation scaling (and MXFP4) therefore fall back to the
unfused AllReduce → RMSNorm → quantize sequence — three kernel launches plus intermediate bf16
round-trips — before every tensor-parallel linear.

Modifications

Adds a per-token FP8 variant (and a 1-stage MXFP4 variant) of the fused kernel, collapsing those
three launches into a single aiter kernel. The change is additive and gated: when the per-token
fast path is not applicable it falls back to the existing per-group / plain AR+RMSNorm path with
no behavior change.

Key changes

Collective entry points (distributed/parallel_state.py, distributed/communication_op.py)

  • tensor_model_parallel_fused_allreduce_rmsnorm_quant_per_token returns
    (fp8_output, residual_out, per_token_scale) with per_token_scale shaped (M, 1), or None
    when the backend cannot fuse (fallback signal).
  • tensor_model_parallel_fused_allreduce_rmsnorm_mxfp4_quant plus the
    _should_use_1stage_mxfp4_ar heuristic for the MXFP4 1-stage path.
  • Device-communicator methods fused_allreduce_rmsnorm_quant_per_token /
    fused_allreduce_rmsnorm_mxfp4_quant, calling aiter
    custom_fused_ar_rms_quant(post_per_token_quant=True).

LayerNorm hook (layers/layernorm.py)

  • forward_with_allreduce_fusion_quant_per_token (plus the shared
    _forward_with_allreduce_fusion_quant_per_token impl), with a cached _aiter_per_token_quant
    functor and a non-aiter / residual is None fallback.

Communicator gating (layers/communicator.py)

  • _try_fused_allreduce_rmsnorm_quant dispatches to the per-token path when
    quant_format == "fp8_per_token" and the LayerNorm exposes the per-token hook; otherwise it
    falls through to per-group / unfused.

Quant fast path (layers/quantization/fp8_utils.py)

  • Consume pre-quantized (fp8, per_token_scale) inputs directly, skipping the linear's internal
    quantization.

Qwen3.5 wiring (models/qwen3_5.py)

  • _detect_fused_ar_quant_format returns "fp8_per_token" under
    SGLANG_USE_AITER_FP8_PER_TOKEN; consumes the fused 2-tuple/3-tuple handoff in both the GDN
    linear-attention (in_proj_qkvz / in_proj_ba) path and the full-attention (qkv_proj) path.

Enablement

SGLANG_USE_AITER=1 + --enable-aiter-allreduce-fusion (existing), plus
SGLANG_USE_AITER_FP8_PER_TOKEN=1 to select the per-token format. MXFP4 checkpoints auto-select
the MXFP4 variant. ROCm/aiter/gfx95-gated; other backends and non-eligible configs are unchanged.

Accuracy Tests

The fused kernel is numerically equivalent to the reference AR → RMSNorm → per-token quant
sequence; when the fast path is unavailable it defers to the existing per-group / plain path.

Verified on Qwen3.5-397B MXFP4-AttnFP8 with GSM8K (1319q): fusion ON = 0.932, and toggling the
per-token FP8 fusion off (SGLANG_DISABLE_FUSED_AR_QUANT) leaves accuracy unchanged at 0.932,
confirming the fallback is numerically faithful.

Unit tests: test/registered/ops/test_aiter_allreduce_fusion_amd.py adds per-token FP8 and MXFP4
cases (fused vs unfused reference) and the None-fallback contract.

Speed Tests and Profiling

Kernel benchmark: benchmark/kernels/all_reduce/benchmark_fused_ar_rms_quant_amd.py adds
per-token FP8 and MXFP4 coverage against the 3-launch baseline.

End-to-end

Setup: Qwen3.5-397B MXFP4-AttnFP8, TP2 on MI355X, fp8 KV cache, --attention-backend aiter,
random dataset, output len 1024, range ratio 0.8, num-prompts = 10 x cc. (1k,1k) uses
--enable-mixed-chunk; (8k,1k) does not. The arms differ only by
SGLANG_USE_AITER_FP8_PER_TOKEN (OFF = per-token fusion disabled → per-group/plain fallback;
ON = per-token fused path). Accuracy is identical between arms (GSM8K 1319q = 0.932 both, see
Accuracy Tests), so this is a like-for-like numerical comparison.

Total token throughput (tok/s), higher is better:

cc (1k,1k) OFF (1k,1k) ON Δ (8k,1k) OFF (8k,1k) ON Δ
4 804.5 805.4 +0.11% 3265.4 3299.7 +1.05%
8 1401.0 1411.1 +0.72% 5357.8 5413.8 +1.05%
16 2129.0 2182.1 +2.49% 7729.3 7749.9 +0.27%
32 3097.5 3168.6 +2.29% 10341.4 10327.2 -0.14%
64 4578.9 4590.5 +0.26% 13255.2 13269.0 +0.10%

Mean TPOT (ms/token), lower is better:

cc (1k,1k) OFF (1k,1k) ON Δ (8k,1k) OFF (8k,1k) ON Δ
4 9.57 9.56 -0.10% 10.38 10.27 -1.06%
8 10.95 10.92 -0.27% 12.58 12.45 -1.03%
16 14.63 14.28 -2.39% 17.58 17.54 -0.23%
32 19.99 19.53 -2.30% 26.22 26.23 +0.04%
64 27.05 26.98 -0.26% 41.44 41.36 -0.19%

The fused path is faster or equal in 9 of 10 cells, and TPOT moves in lockstep with throughput,
which is the signature of removing launches from a decode-bound step rather than of noise. The
largest gains are ~2.4% at (1k,1k) cc=16/32; the rest is neutral. That is the expected magnitude:
the fused region is a small share of a decode step, so collapsing three launches into one removes
real work without dominating end-to-end serving throughput.

Caveat on provenance: these are single samples per cell, and the ON and OFF arms were collected in
separate sessions rather than interleaved, so deltas below ~1% should be read as neutral.

Applicability

The saving is largest when the consumer of the fused output accepts the quantized tensor directly.
If the downstream linear is not quantized in a matching format, the kernel must still produce a
bf16 side output, and the benefit reduces to the removed launches alone.

Checklist

Review and Merge Process

  1. Ping Merge Oncalls to start the process. See the PR Merge Process.
  2. Get approvals from CODEOWNERS and other reviewers.
  3. Trigger CI tests with comments or contact authorized users to do so.
    • Common commands include /tag-and-rerun-ci, /tag-run-ci-label, /rerun-failed-ci
  4. After green CI and required approvals, ask Merge Oncalls or people with Write permission to merge the PR.

CI States

Latest PR Test (Base): ✅ Run #31792406179
Latest PR Test (Extra): ❌ Run #31792406163

@github-actions github-actions Bot added quant LLM Quantization amd labels Jun 30, 2026

@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 support for fused AllReduce + RMSNorm + quantization (including MXFP4, per-group FP8, and per-token FP8) on AMD GPUs using AITER. It adds benchmarking scripts, updates communication operations, integrates the fused paths into layer communicators and layer normalization modules, and updates the Qwen 3.5 model implementation to leverage these optimizations. The feedback highlights several critical issues: the platform check helper _is_gfx95_supported is evaluated as a bare variable instead of being called as a function, which bypasses the check; the benchmark timing function should disable gradient tracking to prevent memory overhead; and potential runtime crashes should be mitigated by handling 3D tensors in _should_use_1stage_mxfp4_ar and adding defensive checks for 1D tensors in _maybe_transpose_aiter_bpreshuffle_scale.

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 thread python/sglang/srt/layers/communicator.py
Comment thread python/sglang/srt/layers/communicator.py
Comment thread benchmark/kernels/all_reduce/benchmark_fused_ar_rms_quant_amd.py
Comment thread python/sglang/srt/distributed/parallel_state.py
Comment thread python/sglang/srt/layers/layernorm.py Outdated
@mqhc2020
mqhc2020 force-pushed the marv/ar_norm_per_token_quant_fusion branch 2 times, most recently from 57713e8 to 2c9514f Compare June 30, 2026 09:29
@mqhc2020
mqhc2020 force-pushed the marv/ar_norm_per_token_quant_fusion branch from 950810f to fd23e99 Compare July 24, 2026 04:22
Conflicts:
	python/sglang/srt/distributed/parallel_state.py
@gemini-code-assist

Copy link
Copy Markdown
Contributor

Caution

The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased.

@mqhc2020

Copy link
Copy Markdown
Contributor Author

@amd-bot ci-status

@amd-bot

amd-bot commented Aug 13, 2026

Copy link
Copy Markdown

@mqhc2020

CI Status for PR #29723

Merge verdict:Not ready to merge. PR CI is incomplete (the AMD ROCm 7.2 run still has a queued mi35x-8-gpu-disaggregation job, and stage-c was fast-fail-skipped) and — most importantly — this PR's core new code is not exercised by any test that ran. The two tests that directly validate the fused per-token AR+RMSNorm+quant path did not run: the dedicated op test is in the skipped stage-c suite, and the qwen3.5 accuracy eval is nightly-only. Every executed failure is either infra (NVIDIA/AMD dependency download) or an unrelated perf/graph-capture flake on non-qwen3.5 models (llama, qwen2_moe, NPU Kimi/MiniMax) — a green-looking signal here would prove nothing about the change.

Caution

This PR's changed code is not verified by the CI that ran.

  • test/registered/ops/test_aiter_allreduce_fusion_amd.py (+276, the direct unit test of the new kernel) is registered to suite stage-c-test-large-8-gpu-amd, which was SKIPPED on this run (stage-c-test-large-8-gpu-amd-rocm720 + -mi35x both skipped) via the fast-fail cascade after stage-b failed.
  • test/registered/amd/accuracy/mi35x/test_qwen35_eval_mi35x.py is registered nightly=True (nightly-amd-accuracy-8-gpu-mi35x-qwen35) → never runs on PR CI.
  • The stage-b-test-large-8-gpu-mi35x-disaggregation-amd-rocm720 job is still queued → AMD ROCm 7.2 run has not finished.
  • Before merge, the author must get stage-c-test-large-8-gpu-amd green (rerun once stage-b infra flakes clear, or add bypass-fastfail to force full signal — sparingly), and run the nightly qwen3.5 MXFP4/FP8 eval manually to confirm accuracy of the new path.

Changed files: distributed/parallel_state.py (+102), layers/communicator.py (+175/-13), layers/layernorm.py (+110), distributed/communication_op.py (+31), layers/quantization/fp8_utils.py (+25), models/qwen3_5.py (+119/-21), models/interns2_mobius.py, 2 test files, 1 benchmark, 1 eval config.

Executed CI failure attribution: AMD: 3 executed test failures (0 clearly related) + 1 infra install-fail · Others: NVIDIA 3 executed (0 related) + 2 fast-fail-skipped · NPU 4 (0 related). AMD ROCm 7.2 run still has 1 job pending — not counted as passed.

AMD Executed Failures

Job Test File Test Function Error Related? Why
stage-b-1gpu-small (0) test/registered/moe/test_torch_compile_moe.py graph capture (qwen2_moe) HIPBLAS_STATUS_INTERNAL_ERROR when calling hipblasCreate during CUDA graph capture 🟢 Crash is in aiter plain gemm_a16w16/hipblasCreate on qwen2_moe — not the PR's qwen3.5 fused path; looks like a runner/driver graph-capture flake
stage-b-2gpu-large (1) test/registered/perf/test_bench_one_batch_2gpu.py throughput assert -1 not greater than 200 (server crashed: hipblasCreate fail during capture on llama.py + aiter_ipc DistStoreError) 🟡 2-GPU TP path is PR-adjacent, but crash is in aiter plain gemm on llama during graph capture, not the PR fused kernel; same HIPBLAS symptom as row 1. Rerun to confirm env flake
stage-b-1gpu-large (1) test/registered/perf/test_bench_serving_1gpu_part2.py median_ttft assert 251.4 not less than 150 🟢 1-GPU (no TP all-reduce fusion). Perf-threshold flake, unrelated to the AR path
stage-b-2gpu-large (0) N/A (install step) N/A Install dependencies step failed (docker manifest sgl-dev:...-rocm720-mi30x-20260812 not found); test never ran 🟢 Infra: image/deps pull failure, not a code failure

Other Executed Failures

Job Test File Test Function Error Related? Why
base-c-4-gpu-h100 (1) N/A (setup) N/A flashinfer_cubin-0.6.17...whl download → http2 refused stream, exit 2 🟢 Network/infra during dep install; AMD-only PR
base-c-8-gpu-h20 (0), (1) N/A (setup) N/A same flashinfer download failure 🟢 Same infra cause
base-c-4-gpu-gb300 (0) (1) N/A N/A Fast-fail: skipping — root cause base-c-4-gpu-h100 (1) 🟢 Cascade skip, not executed
base-c-perf-16-npu-a3 test/registered/npu/performance/kimi_k2_6/test_npu_kimi_k2_6_w4a8_8p_in3k5_out1k5_20ms.py assert_metrics 21.13 not <= 21.0 🟢 NPU perf threshold on Kimi-K2; AMD-only PR
base-c-perf-8-npu-a3 test/registered/npu/performance/minimax_m2_5/test_npu_minimax_m2_5_w8a8_4p_in64k_out1k_prefix90_50ms.py assert_metrics 55.52 not <= 51.0 🟢 NPU perf threshold on MiniMax; unrelated
base-c-perf-2/4-npu-a3 NPU perf command terminated exit 1 🟢 NPU perf; unrelated

Gate/rollup jobs (pr-test-finish, pr-test-npu-finish, wait-for-stage-b-amd-rocm720, pr-test-extra-finish, pr-test-amd-extra-finish, call-gate / pr-gate) are fast-fail/aggregator reflections of the above, not independent failures.

Details / what to do before merge

  1. Verify the changed code actually runs (highest priority). Get stage-c-test-large-8-gpu-amd green so test_aiter_allreduce_fusion_amd.py executes — rerun the AMD ROCm 7.2 workflow after the stage-b infra flakes clear, or add the bypass-fastfail label to force the downstream stage despite the unrelated stage-b failures (uses more CI, use sparingly).
  2. Run the accuracy eval manually. The qwen3.5 MXFP4/FP8 correctness path (test_qwen35_eval_mi35x.py, Qwen3.5-397B-A17B-MXFP4-AttnFP8.yaml) is nightly-only; trigger it (or the nightly-amd-accuracy-8-gpu-mi35x-qwen35 suite / nightly AMD workflow) to confirm the new per-token fused path is numerically correct.
  3. Wait for / rerun the pending mi35x 8-gpu disaggregation job before treating AMD as complete.
  4. Re-check the 🟡 2-GPU crash on rerun. test_bench_one_batch_2gpu crashed via hipblasCreate → HIPBLAS_STATUS_INTERNAL_ERROR during graph capture on llama (not qwen3.5) — the same env symptom as the qwen2_moe failure — so it's most likely a runner/driver flake, but since the PR touches the TP all-reduce/IPC path, confirm it reproduces green on a rerun rather than assuming.
  5. NVIDIA/NPU failures need no action from this PR — pure infra (flashinfer wheel download) and NPU perf thresholds on unrelated models; the author should still confirm required NVIDIA base-c jobs pass on a rerun since they never executed real tests.

Generated by amd-bot using Claude Code CLI

@mqhc2020

Copy link
Copy Markdown
Contributor Author

@amd-bot ci-status

@amd-bot

amd-bot commented Aug 13, 2026

Copy link
Copy Markdown

@mqhc2020

CI Status for PR #29723

Merge verdict:Not ready to merge — CI is still running and incomplete. A new commit (cff2fe5) re-triggered CI ~2 min before this snapshot; NVIDIA is at base-a, AMD ROCm 7.2 is at stage-a, and no downstream stage has been reached. The 4 red checks are not real failures — they are the two "Extra" workflows correctly gating off because the PR lacks the run-ci-extra label. No test that validates this PR's core new code has run yet, so nothing about the fused AR+RMSNorm+quant path is verified.

Caution

This PR's core value (per-token FP8 / MXFP4 fused all-reduce→RMSNorm→quant) is exercised only by two stage-c suites in PR Test ROCm 7.2 (AMD), both currently pending behind stage-a/stage-b (fast-fail gated):

  • stage-c-test-large-8-gpu-amdtest/registered/ops/test_aiter_allreduce_fusion_amd.py (the direct kernel op test), 8-GPU MI300.
  • stage-c-test-large-8-gpu-amd-mi35xtest/registered/amd/perf/mi35x/test_qwen35_fp8_ar_fusion_mi35x.py (qwen3.5 accuracy), 8-GPU MI35x.

The third changed test, test/registered/amd/accuracy/mi35x/test_qwen35_eval_mi35x.py, is registered nightly=True and will not run on PR CI at all. Because stage-c runs only when no upstream AMD job fast-fails, an unrelated stage-a/stage-b flake would silently skip both coverage suites (no bypass-fastfail label set). A green run cannot be claimed until those two stage-c suites actually execute and pass.

Changed source: distributed/parallel_state.py (+102), distributed/communication_op.py (+31), layers/communicator.py (+175/-13), layers/layernorm.py (+110), layers/quantization/fp8_utils.py (+25), models/qwen3_5.py (+119/-21), models/interns2_mobius.py (+4/-3). Plus 3 test files + 1 lm_eval config.

Executed CI failure attribution: AMD: 0 real failures (main gate still running) · Others: 0 real failures (all "Extra" reds are benign label opt-outs). Real test jobs still pending/in-progress — do not read as passed.

Benign gate "failures" (not caused by this PR)

Job Reason Related?
call-gate / pr-gate (PR Test Extra) Missing required label 'run-ci-extra' → exit 1 by design; all extra jobs skipped 🟢 No
pr-test-extra-finish Rollup reports the gated-off extra run as failed 🟢 No
call-gate / pr-gate (PR Test Extra AMD) Same missing run-ci-extra label 🟢 No
pr-test-amd-extra-finish Rollup of the gated-off AMD extra run 🟢 No

Still running (decisive signal)

  • PR Test (NVIDIA, run 31670430117) — base-a-test-cpu / base-a-test-1-gpu-small in progress; base-b/base-c not reached.
  • PR Test ROCm 7.2 (AMD) (run 31670430136) — stage-a-test-1-gpu-small-amd-rocm720 in progress; the two stage-c coverage suites pending.
  • Also pending: PR Test (NPU / XPU / Xeon / Arm64 / MLX). MUSA + sgl-router already ✅.

What to do before merge

  • Wait for PR Test ROCm 7.2 (AMD) to reach and pass stage-c-test-large-8-gpu-amd-rocm720 and stage-c-test-large-8-gpu-amd-mi35x-rocm720 — these are the only PR-CI signals for the new fused path. If either is skipped due to an upstream fast-fail, the core change is untested regardless of overall green.
  • To validate the nightly-only accuracy path (test_qwen35_eval_mi35x.py) before merge, trigger it manually (its suite nightly-amd-accuracy-8-gpu-mi35x-qwen35 does not run on PRs).
  • Since changes to layernorm.py / communicator.py / fp8_utils.py are shared (non-AMD) files, confirm the NVIDIA base-*/stage-* and NPU/XPU runs finish green too — the new code is aiter-gated, but the fallback branches touch common paths.

Note: a prior ci-status comment (2026-08-13 03:08 UTC) was for an earlier commit where stage-c had already been fast-fail-skipped; the branch was updated since, so CI is now back at the beginning.

Generated by amd-bot using Claude Code CLI

Conflicts:
	python/sglang/srt/models/qwen3_5.py
@mqhc2020

Copy link
Copy Markdown
Contributor Author

@amd-bot ci-status

@amd-bot

amd-bot commented Aug 13, 2026

Copy link
Copy Markdown

@mqhc2020

CI Status for PR #29723

Merge verdict: 🚫 Not ready — PR CI is incomplete, and the changed code has not been exercised yet. The only completed "failures" are by-design gate cascades in the Extra workflows (missing run-ci-extra label), not code failures. The real PR-gating pipelines — NVIDIA pr-test.yml (queued) and AMD pr-test-amd-rocm720.yml (in-progress, still at stage-a) — have not concluded. Do not read this as green.

Caution

This PR's core value — the AMD aiter fused AllReduce→RMSNorm→per-token-FP8/MXFP4 quant kernels — is not yet verified by any completed PR-CI test:

  • test/registered/ops/test_aiter_allreduce_fusion_amd.py (suite stage-c-test-large-8-gpu-amd) and test/registered/amd/perf/mi35x/test_qwen35_fp8_ar_fusion_mi35x.py (suite stage-c-test-large-8-gpu-amd-mi35x, runs on linux-mi35x-gpu-8) run only at stage-c of the AMD ROCm 7.2 pipeline, which is gated behind stage-b and has not run yet (AMD is still executing stage-a).
  • test/registered/amd/accuracy/mi35x/test_qwen35_eval_mi35x.py is registered nightly=True → it never runs on PR CI. The GSM8K accuracy of the new MXFP4/FP8 path must be confirmed via a nightly run or manual invocation before merge.
    Green PR CI will only be meaningful once the stage-c AMD 8-GPU jobs (regular + mi35x) actually conclude.

Changed files: mostly AMD paths — distributed/parallel_state.py (+102), layers/communicator.py (+175/-13), layers/layernorm.py (+110), models/qwen3_5.py (+118/-21), distributed/communication_op.py (+31), layers/quantization/fp8_utils.py (+25), plus 3 AMD/mi35x tests and a kernel benchmark.

Executed CI failure attribution: AMD: 2 failures (0 related) · Others: 2 failures (0 related). All 4 are Extra-workflow gate cascades. Main NVIDIA + AMD + NPU/XPU/Xeon/Arm64/MLX pipelines are still pending — not counted as passed.

Other Executed Failures

Job Test File Test Function Error Related? Why
call-gate / pr-gate N/A N/A Missing required label 'run-ci-extra' → exit 1 🟢 By-design opt-in gate; PR has no run-ci-extra label. Not a code failure.
pr-test-extra-finish N/A N/A Fails because gate failed (cascade) 🟢 Aggregate/finish job; all extra tests skipped.

AMD Executed Failures

Job Test File Test Function Error Related? Why
call-gate / pr-gate N/A N/A Missing required label 'run-ci-extra' → exit 1 🟢 Same opt-in gate for Extra (AMD) workflow.
pr-test-amd-extra-finish N/A N/A Fails because gate failed (cascade) 🟢 Aggregate/finish job; all extra-AMD tests skipped.

Details / what to do before merge

  • Wait for the real pipelines to finish. Re-check when NVIDIA pr-test.yml and AMD pr-test-amd-rocm720.yml reach stage-c and conclude. Green MUSA/sgl-router alone do not cover this PR.
  • Confirm the AMD stage-c jobs actually run and pass, specifically stage-c-test-large-8-gpu-amd-rocm720 (runs test_aiter_allreduce_fusion_amd.py) and stage-c-test-large-8-gpu-amd-mi35x-rocm720 on linux-mi35x-gpu-8 (runs test_qwen35_fp8_ar_fusion_mi35x.py). These are the tests that exercise the new fused kernel and its fallback path.
  • The mi35x GSM8K eval (test_qwen35_eval_mi35x.py) is nightly-only — trigger it via nightly or manual dispatch to verify accuracy of the new MXFP4/per-token-FP8 path before merge, since PR CI will never run it.
  • The Extra-workflow failures are safe to ignore for merge, or add the run-ci-extra label (e.g. /tag-and-rerun-ci extra) if you want the 4-GPU/8-GPU extra coverage to actually run (consumes more CI).

Generated by amd-bot using Claude Code CLI

@sogalin sogalin left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Looks good to me now, we cover test cases.

@kkHuang-amd

Copy link
Copy Markdown
Collaborator

I found several blockers on the current head 4def6486:

  1. The decisive MI35X Qwen3.5 job still fails in this PR's code. During CUDA-graph capture, _select_fused_ar_input_for_linear(..., in_proj_ba) raises TypeError: MergedColumnParallelLinear cannot consume fused AR quant tuple input. If a consumer cannot accept the exact emitted quant format, the producer must emit a BF16 sidecar and the consumer must select it; a cached boolean “accepts any tuple” is not a sufficient producer/consumer format contract. Please add a regression using the failing Qwen3.5 merged in_proj_ba path.

  2. Arbitrary collective exceptions cannot safely become local fallback. The fused collective wrappers catch every exception and return None. Once any rank enters the collective, ranks can diverge into different fallback behavior, hang, double-reduce, or reuse partially modified communicator state. Return None only from deterministic preflight checks before collective entry; after launch, propagate the failure or use a rank-consistent failure protocol.

  3. The new quantized methods omit the existing piecewise-CUDA-graph guard. In the AITER _IS_CAPTURING/non-capturing-stream state, the custom wrapper can return dummy zero outputs; SGLang currently treats a non-None tuple as real data. Mirror the unquantized guard and add capture/replay correctness tests for per-token FP8 and MXFP4.

  4. The one-stage heuristic can violate the documented 80-token kernel limit. total_bytes <= 128 KiB alone accepts cases such as BF16 [128,512]. Require token_num <= 80 in addition to the measured hidden-size crossover.

Additional issues:

  • The prequantized FP8 tuple path loses the original activation dtype and hardcodes BF16; carry dtype and test FP16.
  • InternS2 computes capability but does not pass the required quant_format/fuse_quant wiring, so its fusion is ineffective.
  • The current Qwen per-token test uses the all-MXFP4 checkpoint rather than MXFP4-AttnFP8; the per-token variant never ran after the earlier failure.
  • The direct fused-op test is registered nightly=True in a PR stage-C suite, so it did not execute in current-head CI.

The branch is also ~176 main commits behind and overlaps semantically with #34502 in five shared infrastructure files. Please rebase and reconcile a single fused-AR/tuple API before merge.

Finally, the E2E benchmark toggle is confounded: SGLANG_USE_AITER_FP8_PER_TOKEN changes both quant format and fusion. Keep per-token quant enabled in both arms and toggle only SGLANG_DISABLE_FUSED_AR_QUANT, with interleaved repeated runs and raw results.

@kkHuang-amd kkHuang-amd left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Inline findings for current head 4def6486, expanding the summary in #29723 (comment). Each comment identifies the concrete changed line, reachable impact, and requested fix.

if len(hidden_states) == 2 and _linear_accepts_fp8_tuple(linear):
if len(hidden_states) == 2 and _linear_accepts_quant_tuple(linear):
return hidden_states
raise TypeError(

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

[P0] The current-head MI35X production test still reaches this raise. During Qwen3.5 CUDA-graph capture the fused producer returns a 2-tuple, while merged in_proj_ba cannot consume that exact format and has no BF16 sidecar to select. The cached boolean capability therefore does not keep producer and both consumers in agreement. Please define an explicit emitted-format/consumer-format contract and always request/select a BF16 sidecar when either GDN projection cannot consume the exact quant tuple. Add a regression using the failing Qwen3.5-397B-A17B-FP8 merged in_proj_ba path.

if hidden_size == 7168:
# CUDA-graph microbench: direct MXFP4 epilogue is faster through 56
# tokens, while fallback wins from 64 tokens onward.
return tokens <= 56

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

[P0] This can select the 1-stage kernel above its documented 80-token hard limit. For example BF16 [128, 512] is exactly 128 KiB and returns true here despite having 128 tokens. Please require tokens <= 80 for every default 1-stage decision; keep the measured K=7168 cutoff as an additional restriction, not a replacement for the hard limit. Add boundary tests for 80/81 tokens.

use_1stage_ar = _should_use_1stage_mxfp4_ar(input_)

try:
return ca_comm.custom_fused_ar_rms_mxfp4_quant(

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

[P0] The new quantized collectives omit the existing TC-piecewise CUDA-graph guard. AITER's custom communicator can return dummy zero outputs when its global capture state is active but the current stream is not capturing; this non-None tuple is then treated as real activations/residuals/scales. Mirror fused_allreduce_rmsnorm's capture-state handling before calling the MXFP4 and per-token wrappers, and add piecewise capture/replay correctness tests for both formats.

use_1stage_ar,
emit_bf16=emit_bf16,
)
except Exception:

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

[P0] An arbitrary collective runtime failure cannot safely become a per-rank local fallback. Once one or more ranks have entered the fused collective, independently returning None can make peers hang, run a second all-reduce, or consume partially mutated communicator/tensor state. Return None only from deterministic preflight checks before collective entry. Once the backend call begins, propagate the exception or implement a rank-consistent failure protocol. The same issue exists in the per-token wrapper's broad catch.

WQ=weight.T,
x_scale=x_scale,
w_scale=weight_scale,
dtype=torch.bfloat16,

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

[P1] This tuple path loses the producer's original activation dtype. The normal path emits input.dtype, but a prequantized FP16 activation is silently promoted to BF16 here. Carry the intended output dtype in the tuple/API (as the static prequantized path does), pass it to the GEMM, and add FP16 plus BF16 tuple-path coverage.

enable_fused_ar_quant = (
_enable_qwen35_fused_ar_quant()
and _linear_accepts_fp8_tuple(self.linear_attn.in_proj_qkvz)
and _linear_accepts_quant_tuple(self.linear_attn.in_proj_qkvz)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

[P1] InternS2 computes capability but never completes the new communicator wiring. Both decoder forwards still call prepare_attn_and_capture_last_layer_outputs without the required quant_format, emit_bf16, and fuse_quant=True, so the communicator does not select the per-token/MXFP4 path (and the prior opt-in is effectively disabled). Please store the consumer format and pass the full fusion contract from both forwards, or remove this ineffective opt-in.

register_amd_ci(
est_time=480,
suite="stage-c-test-large-8-gpu-amd-mi35x",
nightly=True,

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

[P1] This registration makes the direct fused-op coverage effectively orphaned. PR stage-C invokes the suite without --nightly, while no nightly workflow invokes this exact suite; current-head logs contain no execution of this file. Either remove nightly=True for the PR suite or register a real nightly suite that is actually dispatched, then provide a non-skipped current-head run.

)
QWEN35_MXFP4_MODEL_PATH = os.environ.get(
"QWEN35_MXFP4_MODEL_PATH",
"amd/Qwen3.5-397B-A17B-MXFP4",

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

[P1] This all-MXFP4 checkpoint does not establish the claimed per-token FP8 path. Setting SGLANG_USE_AITER_FP8_PER_TOKEN=1 cannot create FP8 attention consumers when the checkpoint itself is all MXFP4, and current CI aborts in the preceding per-group case before this method starts. Use the MXFP4-AttnFP8 checkpoint, assert that both per-token-FP8 and MXFP4 epilogues actually fire, and calibrate the acceptance threshold against that checkpoint's baseline.

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.

9 participants