Skip to content

[AMD] support gfx1250 on ROCM 10 - #36871

Merged
HaiShaw merged 101 commits into
mainfrom
amd_helios
Aug 31, 2026
Merged

HaiShaw merged 101 commits into
mainfrom
amd_helios

Conversation

@yctseng0211

@yctseng0211 yctseng0211 commented Aug 28, 2026 •

Copy link
Copy Markdown
Collaborator

Motivation

Based on #32754 (gfx1250 enablement) and #36434 (ROCm 10 release images), see those PRs for the details of each.

On top of them, this adds the gfx1250-rocm1000 Docker target on the ROCm 10.0.0 GA wheel channel, gfx1250 detection in the AMD CI dependency installer, and the gfx1250 kernel/model fixes with the MI45x accuracy tests.

Modifications

Accuracy Tests

Verifying on MI455x.

GPT-OSS-120b

HIP_FORCE_DEV_KERNARG=0 ENABLE_CK=0 python3 -m sglang.launch_server   \
--model-path /dockerx/data/models/gpt-oss-120b-w-mxfp4-a-fp8   --tp 1   \
--trust-remote-code   --max-running-requests 128   --mem-fraction-static 0.9   \
--attention-backend triton   --disable-radix-cache   --disable-cuda-graph
1319/1319 [04:20<00:00,  5.07it/s]
Accuracy: 0.845
Invalid: 0.010
Latency: 260.197 s
Output throughput: 1630.529 token/s

DSR1

HIP_FORCE_DEV_KERNARG=0  \
ENABLE_CK=0 SGLANG_USE_AITER=1 AITER_FORCE_A8W4=1 \
AITER_GROUPED_FORCE_SPLIT_K1=1 SGLANG_MOE_SHUFFLE_GFX1250=1 \
ROCM_QUICK_REDUCE_QUANTIZATION=NONE SGLANG_AITER_FP8_PREFILL_ATTN=0 \
SGLANG_AITER_MLA_PERSIST=0 SGLANG_INT4_WEIGHT=0 SGLANG_MOE_PADDING=1 \
SGLANG_SET_CPU_AFFINITY=1 SGLANG_ROCM_FUSED_DECODE_MLA=0 SGLANG_USE_ROCM700A=1 \
AITER_GROUPED_CONTIGUOUS_TOKEN_THRESHOLD=16 \
python3 -m sglang.launch_server \
  --model-path /dockerx/data/models/DeepSeek-R1-0528-MXFP4/ \
  --tensor-parallel-size 1 \
  --trust-remote-code \
  --host 0.0.0.0 \
  --port 8000 \
  --mem-fraction-static 0.90 \
  --chunked-prefill-size 16384 \
  --attention-backend triton \
  --max-running-requests 32 \
  --kv-cache-dtype auto \
  --page-size 64
1319/1319 [11:00<00:00,  2.00it/s]
Accuracy: 0.951
Invalid: 0.000
Latency: 660.142 s
Output throughput: 201.642 token/s

DSv4-Flash

HIP_FORCE_DEV_KERNARG=0 SGLANG_DEFAULT_THINKING=1 \
SGLANG_DSV4_REASONING_EFFORT=max SGLANG_USE_ROCM700A=0 \
SGLANG_HACK_FLASHMLA_BACKEND=unified_kv_triton AITER_BF16_FP8_MOE_BOUND=0 \
AITER_FORCE_A8W4=1 SGLANG_USE_AITER_MOE_GU_ITLV=1 \
SGLANG_OPT_FUSE_MHC_POST_PRE=0 ENABLE_CK=0 SGLANG_USE_AITER=1 \
python3 -m sglang.launch_server --model-path /dockerx/data/models/DeepSeek-V4-Flash \
  --trust-remote-code --tp 1 --attention-backend dsv4 --page-size 256 \
  --mem-fraction-static 0.60 --swa-full-tokens-ratio 0.15 --disable-shared-experts-fusion \
  --tool-call-parser deepseekv4 --reasoning-parser deepseek-v4 --chunked-prefill-size 8192 \
  --cuda-graph-max-bs 256 --max-running-requests 256 --disable-radix-cache \
  --kv-cache-dtype fp8_e4m3
1319/1319 [01:57<00:00, 11.27it/s]
Accuracy: 0.929
Invalid: 0.000
Latency: 117.025 s
Output throughput: 1013.394 token/s

Wan2.2

export SGLANG_DIFFUSION_PLATFORM_OVERRIDE=rocm
export TORCHINDUCTOR_MAX_AUTOTUNE=0
export TORCHINDUCTOR_MAX_AUTOTUNE_POINTWISE=0
export ENABLE_CK=0
export MODEL=/dockerx/data/models/Wan2.2-T2V-A14B-Diffusers/
export PROMPT="A cat and a dog baking a cake together in a kitchen. \
The cat is carefully measuring flour, while the dog is stirring the batter with a wooden spoon. \
The kitchen is cozy, with sunlight streaming through the window."
export SGLANG_DIFFUSION_VAE_CHANNELS_LAST_3D=1
PYTORCH_NO_HIP_MEMORY_CACHING=1 HIP_FORCE_DEV_KERNARG=0 SGLANG_USE_ROCM_VAE_CONV2D_BF16=1 SGLANG_USE_ROCM_FLYDSL=0 SGLANG_DIFFUSION_AITER_FP8_ATTN=0 sglang generate   --model-path "$MODEL"   --log-level info   --prompt "$PROMPT"   --negative-prompt " "   --height 1280   --width 720   --num-inference-steps 8   --num-frames 33   --seed 42   --save-output   --num-gpus 1   --vae-decode-precision fp32   --dit-layerwise-offload true   --dit-cpu-offload false   --vae-cpu-offload false   --text-encoder-cpu-offload false   --enable-torch-compile false

Flux.2-dev

HIP_FORCE_DEV_KERNARG=0 SGLANG_DIFFUSION_PLATFORM_OVERRIDE=rocm ENABLE_CK=0 sglang generate   --model-path /dockerx/data/models/FLUX.2-dev   --prompt "A cinematic shot of a cute cat wearing a tiny astronaut helmet."   --num-gpus 1   --tp-size 1   --save-output --width 1024 --height 1024

Speed Tests and Profiling

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 #33315585657
Latest PR Test (Extra): ❌ Run #33315585527
Latest PR Test (AMD ROCm 7.2): ❌ Run #33315585645

HaiShaw and others added 30 commits July 22, 2026 03:01
…SR1, GPTOSS) (#32043)

Co-authored-by: Kao <akao@amd.com>
Co-authored-by: wunhuang <wunhuang@amd.com>
Co-authored-by: Thomas Wang <1am9trash@gmail.com>
Co-authored-by: Xinyi Song <86638975+RolaoDenthu@users.noreply.github.com>
Co-authored-by: Lin, Soga <soga.lin@amd.com>
Co-authored-by: wunhuang <wunhuang@amd.com>
…4, and GPT-OSS W4A8 MXFP4-FP8

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
rocm700/rocm720 (legacy SDK) require compressed-tensors==0.15.0 while
rocm7_15+ (AMD rock SDK) requires ==0.16.0.  Remove the dep from the
shared runtime_common extra (which caused ResolutionImpossible when a
platform extra tried to override it) and introduce two new explicit
extras instead:

  rocm_legacy = compressed-tensors==0.15.0  (ROCm 7.0 / 7.2)
  rocm_rock   = compressed-tensors==0.16.0  (ROCm 7.15+, rock SDK)

Non-HIP platforms that also consume runtime_common (HPU, MUSA, MPS)
now declare compressed-tensors==0.15.0 directly in their own srt_*
extras so they are unaffected.

docker/rocm.Dockerfile: replace the broad sed relaxation with a
case branch on GPU_ARCH that injects rocm_rock for gfx1250-rocm7_15
and rocm_legacy for all other arches.

scripts/ci/amd/amd_ci_install_dependency.sh: detect the ROCm version
inside the CI container (same pattern used by AITER/MORI blocks) and
append the correct extra to EXTRAS before the sglang pip install.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
…MD platforms

MLX/MPS (Apple Silicon), HPU, and MUSA platforms may not have a wheel
for compressed-tensors==0.15.0, causing ResolutionImpossible.  Restore
the bare (unpinned) entry in runtime_common so those platforms let pip
pick whatever version is available.

The AMD-specific pin is still enforced: rocm_legacy (==0.15.0) and
rocm_rock (==0.16.0) add a precise constraint on top of the bare dep,
which pip satisfies unambiguously for AMD builds.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Add __builtin_trap() + return to the non-void stub so that calling it
on gfx1250 (where the real buffer load intrinsic doesn't exist) fails
loudly instead of silently invoking undefined behavior.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Commit 9579fa3 fixed a gfx1250 hardware bug where tl.dot(fp8, fp8)
returns garbage (~1e34+) for contraction dim K>=128 by upcasting K/V to
the query dtype.  That fix landed without a platform guard, so CUDA and
non-gfx1250 ROCm paths were also affected, changing their fp8 matmul
precision.

Add IS_GFX1250: tl.constexpr (default False) to both _fwd_kernel and
_fwd_kernel_unified.  Gate all seven affected dot sites behind the
constexpr so:
  - gfx1250: upcast K/V to q.dtype (original workaround preserved)
  - all other platforms: restore q.to(k.dtype) / p.to(v.dtype) (original
    pre-workaround behaviour)

Pass IS_GFX1250=_is_gfx1250 from both Python launchers.  Triton compiles
a separate specialisation per constexpr value, so there is no runtime
branch overhead on either platform.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Commit 9579fa3 worked around a gfx1250 hardware bug where tl.dot(fp8, fp8)
returns garbage (~1e34+) for contraction dim K>=128, by keeping q in bf16 and
upcasting K to q.dtype, and by keeping the softmax weights p in fp32 for the
P/V accumulation.  9626ada added an IS_GFX1250 guard for extend_attention.py
but left decode_attention.py unguarded, so CUDA and non-gfx1250 ROCm targets
still paid for the workaround -- most notably the forced fp32 P/V dot in
_fwd_grouped_kernel_stage1.

Add IS_GFX1250: tl.constexpr (default False) to _fwd_grouped_kernel_stage1 and
gate the three affected sites so:
  - gfx1250: workaround preserved
  - all other platforms: restore q.to(K_Buffer.dtype.element_ty), tl.dot(q_k, k)
    and tl.dot(p.to(v.dtype), v)

Pass IS_GFX1250=_is_gfx1250 from _decode_grouped_att_m_fwd.  Triton compiles a
separate specialisation per constexpr value, so neither platform pays a runtime
branch.

Also mark every IS_GFX1250 guard in decode_attention.py and extend_attention.py
with a TODO, so these branches are recognised as temporary and removed once the
underlying gfx1250 issues are resolved.

Verified on gfx1250 that the IS_GFX1250=False path is bitwise identical to the
pre-workaround kernel across 12 configs (bf16 and fp8 KV, including the MLA
576/512 shapes), and that the IS_GFX1250=True path still avoids the fp8 NaN.
@yctseng0211

Copy link
Copy Markdown
Collaborator Author

/rerun-failed-ci

# Conflicts:
#	python/sglang/kernels/ops/attention/dsv4/unified_kv_kernels/paged_decode.py
@yctseng0211

Copy link
Copy Markdown
Collaborator Author

resolved conflict :
python/sglang/kernels/ops/attention/dsv4/unified_kv_kernels/paged_decode.py

yctseng0211 and others added 4 commits August 30, 2026 00:32
The MI45x bring-up ran on upstream Triton 76940ad, carried over from the ROCm
7.14 flavor this image replaced, and it fixes an issue the SDK's Triton still
has. The revision was already recorded on the stage but nothing consumed it:
BUILD_TRITON was 0, and the source-build arm went away with the 7.14 flavor.

Scoped to gfx1250-rocm1000 by exact match. gfx942 and gfx950 keep the Triton
their SDK ships, and the ROCm 7.2 images keep their pinned wheels.
@yctseng0211

Copy link
Copy Markdown
Collaborator Author

PR-base (CUDA) failures are pre-existing on main

Failing test Job This PR main
test_fp8_utils.py base-b-test-1-gpu-large fail fail
test_deepseek_v4.py base-b-test-1-gpu-large fail fail
test_deepseek_v4.py base-b-test-4-gpu-b200 fail fail

The same three reproduce on main in
run 33308180355
(2026-08-30 11:08 UTC), with identical test files and — for test_fp8_utils.py —
a byte-identical RuntimeError once timestamps are stripped.

main was green through
33280258528
(08-29 23:08) and red by 08-30 11:08. Inside that window only one commit touched
layers/quantization/fp8.py and layers/attention/dsv4/indexer.py:

7e751153eb [Config] Round 5.1: the published-side readers ask the bags, and a
           platform fact gets one address (#37086)

That PR routes is_sm90 / is_sm100 / is_sm120 / is_hip through
get_platform(), which lines up with both failing tests. Not from this PR — this
one only touches AMD/ROCm paths.

@HaiShaw HaiShaw 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.

LGTM, gfx specific.

@HaiShaw

HaiShaw commented Aug 31, 2026

Copy link
Copy Markdown
Collaborator

gfx specific changes, failed PR Test Base were known issues.

@HaiShaw
HaiShaw merged commit 3865efc into main Aug 31, 2026
145 of 204 checks passed
@HaiShaw
HaiShaw deleted the amd_helios branch August 31, 2026 08:19
@yctseng0211

Copy link
Copy Markdown
Collaborator Author

Note: origin of the mem_fraction_static aiter guard, and its overlap with #36656

Recording this for traceability, since #36656 has independently surfaced the same
failure mode today.

What this PR carried forward

This PR did not originate the mem_fraction_static guard in
python/sglang/srt/arg_groups/attention_hook.py. It came from
#32754 ("[AMD] Enable gfx1250
Support", closed unmerged), where it lived in server_args.py:

user_set = (getattr(self, "_raw_input", None) or {}).get("mem_fraction_static") is not None
if user_set:
    logger.warning(...)
else:
    self._declare("_handle_attention_backend_compatibility",
                  mem_fraction_static=self.mem_fraction_static * 0.85)

The block moved to arg_groups/attention_hook.py via the config-resolution
refactor (#35905/#35906/#36789/#36896). The version merged here is functionally
equivalent to the #32754 original — same _raw_input predicate, same warn/else
split — with a longer explanatory comment.

It was flagged on #32754, twice, with a confirmed mechanism

@bingxche identified this as a regression on the originating PR:

nightly-4-gpu-mi35x-minimax-m3-rocm720 failed with HIP out of memory. Tried to allocate 67.41 GiB on a 288 GB MI355. test_minimax_m3_tp4_eval_mi35x.py launches
with --attention-backend aiter and an explicit --mem-fraction-static 0.80
(lines 92-100), so the guard takes the effective fraction from 0.80 * 0.85 = 0.68
to 0.80 — roughly 34 GB more static budget. Reproduced identically on cc42c15
and 852a411, same day and runner pool, with main green.

From that review: "This change is not arch-gated and not gfx1250-specific: every
--attention-backend aiter deployment that sets --mem-fraction-static explicitly
with context_len > 8192 silently gets a larger static budget than before."

Same failure mode as #36656

#36656 deletes the * 0.85
block unconditionally. This PR skips it only when the user set the value explicitly.
Different scope, identical behaviour on the intersection — and the intersection is
not empty:

Test Explicit value Before #36656 This PR
test_minimax_m3_tp4_eval_mi35x.py 0.80 0.68 0.80 → OOM 0.80 → OOM
test_streaming_session.py 0.70 0.595 0.70 → OOM 0.70
test_streaming_session_extra.py 0.70 0.595 0.70 → OOM 0.70

The MiniMax row is the same test, same numbers and same signature reported in both
the #32754 review above and the #36656 test
summary

(MF 0.80, 67.32 GiB requested / ~47.4 GiB free, baseline passing at old effective
0.68). The two test_streaming_session* files use
meta-llama/Llama-3.1-8B-Instruct (131072 context, over the 8192 threshold) and
were listed in #36656 as direct AITER OOMs at 16.50 GiB requested; both set
--mem-fraction-static 0.7, so they fall inside this PR's guard as well.

Of the eleven individually named test files in the #36656 report, the other nine do
not set mem_fraction_static explicitly, so they still receive the * 0.85 scaling
under this PR and are unaffected by it; I have not checked the bench-serving,
pipeline-parallel or MORI/NIXL entries. Repo-wide, 427 test files pass
--mem-fraction-static; the exposed subset is those that also run on AMD with aiter
and context_len > 8192.

This PR's AMD CI did not exercise it

The last AMD run on this PR
(33315585645, sha
52a3d8390) failed during the build, on the quick_all_reduce.cuh:512
v_cvt_pk_f16_f32 "instruction not supported on this GPU" error from #34484 — the
runtime tests never started. That build fix landed separately in
#37132, not here. So the earlier
"no new regressions" comparison on this PR does not cover the guard's runtime
behaviour, and there is no post-merge AMD run on main yet.

Open question

Both rationales hold. The guard exists because scaling down a user-chosen value can
push the static budget below the model-weight footprint on a nearly full GPU and
break KV-cache allocation outright. The #32754 and #36656 data show the 0.85 reserve
is load-bearing for the AITER workspace. A floor-based rule — keep the * 0.85 and
skip it only when the scaled result would fall below the model-weight footprint —
would satisfy both, rather than keying on whether the value was set explicitly.

Waiting on the next AMD nightly to see whether the predicted OOMs materialise on
main before proposing a change.

StevenChenSE pushed a commit to StevenChenSE/sglang that referenced this pull request Sep 6, 2026
Co-authored-by: HAI <hixiao@gmail.com>
Co-authored-by: Kao <akao@amd.com>
Co-authored-by: wunhuang <wunhuang@amd.com>
Co-authored-by: Thomas Wang <1am9trash@gmail.com>
Co-authored-by: Xinyi Song <86638975+RolaoDenthu@users.noreply.github.com>
Co-authored-by: Lin, Soga <soga.lin@amd.com>
Co-authored-by: kk <43161300+kkHuang-amd@users.noreply.github.com>
Co-authored-by: Bingxu Chen <bingxche@amd.com>
Co-authored-by: sogalin_codegen <39478626+sogalin@users.noreply.github.com>
Co-authored-by: Thomas Wang <thomawan@amd.com>

This branch was previously deployed

1 inactive deployment
prod — 52a3d839 Deployed Aug 30, 2026 by yctseng0211 via publish-gfx1250 #10
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

amd bypass-fastfail deepseek dependencies Pull requests that update a dependency file high priority jit-kernel run-ci CI: run the baseline test suite on this PR sgl-kernel

Projects

None yet

Development

Successfully merging this pull request may close these issues.

8 participants