Repository navigation
[AMD] support gfx1250 on ROCM 10 - #36871
Conversation
…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.
|
/rerun-failed-ci |
# Conflicts: # python/sglang/kernels/ops/attention/dsv4/unified_kv_kernels/paged_decode.py
|
resolved conflict : |
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.
PR-base (CUDA) failures are pre-existing on main
The same three reproduce on main in main was green through That PR routes |
|
gfx specific changes, failed PR Test Base were known issues. |
Note: origin of the
|
| 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.
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>
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
DSR1
DSv4-Flash
Wan2.2
Flux.2-dev
Speed Tests and Profiling
Checklist
Review and Merge Process
/tag-and-rerun-ci,/tag-run-ci-label,/rerun-failed-ciCI States
Latest PR Test (Base): ❌ Run #33315585657
Latest PR Test (Extra): ❌ Run #33315585527
Latest PR Test (AMD ROCm 7.2): ❌ Run #33315585645