[release/2.13] [ROCm] Add initial support for gfx1250 + disable CK GEMM when gfx1250-only - #3507
Draft
ethanwee1 wants to merge 2 commits into
Draft
[release/2.13] [ROCm] Add initial support for gfx1250 + disable CK GEMM when gfx1250-only#3507ethanwee1 wants to merge 2 commits into
ethanwee1 wants to merge 2 commits into
Conversation
#3421) ## Summary Validation theRock Run: https://github.com/ROCm/TheRock/actions/runs/29045360944 Cherry-pick of upstream [pytorch#188597](pytorch#188597) onto `release/2.12` to add initial ROCm support for **gfx1250 (CDNA5)**. Co-authored-by: @glen-amd ### What this PR enables - Gates gfx1250-specific behavior on ROCm 7.14+ (the support floor), applied consistently across CUDABlas.cpp, ScaledBlas.cpp, CUDAHooks.cpp, and the test helpers. - Adds gfx1250 to the hipBLASLt preferred/supported arch lists and the hipSparseLt support check (ROCm 7.14+). - Extends the scaled GEMM / MX-format paths (block-wise Float8_e8m0fnu, mxfp8, mxfp4) to gfx1250. - Sets the gfx1250 shared-memory limit to 320 KB (vs 160 KB on gfx950), matching on the suffix-stripped arch name. - Adds CDNA2/CDNA3/CDNA5 "or later" arch helpers and updates FP8 / MX GEMM test gating in common_cuda.py. - Builds Composable Kernel (CK) GEMM as a separate ck_gemm library with gfx1250 filtered out of HIP_ARCHITECTURES, so a multi-arch build that includes gfx1250 keeps CK GEMM for the other archs instead of failing to compile for gfx1250 (composable_kernel has no gfx1250 support yet). The filter is removed once CK supports gfx1250. - Extends the gfx942 nontemporal vectorized-load path in MemoryAccess.cuh to gfx1250. - Refactors duplicated arch-check logic in CUDABlas.cpp, ScaledBlas.cpp, and cuSPARSELtOps.cpp per review feedback. - Writes the CublasHandlePool workspace archs in full (gfx942/gfx950/gfx1250). Enabled via the main merge: - Flash attention and memory-efficient attention on gfx1250 via AOTriton 0.12.1b (from pytorch#188242), which this branch merges in. What is not enabled yet: - CK SDPA on gfx1250: composable_kernel has no gfx1250 support; the CK SDPA target auto-filters to gfx942/gfx950. - CK GEMM on gfx1250: filtered out (see above) until composable_kernel supports it. - FP8 grouped GEMM on gfx1250: MSLK builds only gfx942/gfx950. - int4 mm on gfx1250: the tinygemm MFMA kernel needs a WMMA port; the ops return a clear "not supported yet" error and the int4 unit tests skip gfx1250. CI note: - gfx1250 is not added to the ROCm 7.2 (rocm-n) docker image, whose HIP compiler cannot target gfx1250. gfx1250 CI belongs on the nightly ROCm image and is added in a follow-up. Co-authored-by: @glen-amd ## References - Upstream PR: pytorch#188597 - Related AOTriton bump: pytorch#188242 --------- Co-authored-by: Prachi Gupta <prachi.gupta@amd.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Co-authored-by: Jeff Daily <jeff.daily@amd.com>
composable_kernel has no gfx1250 support yet. When gfx1250 is the only arch in PYTORCH_ROCM_ARCH, disable USE_ROCM_CK_GEMM before the -DUSE_ROCM_CK_GEMM flag is added in cmake/Dependencies.cmake, so both the CK GEMM sources and their call sites are compiled out consistently and single-arch gfx1250 builds link cleanly. Multi-arch builds are unaffected (ck_gemm is still built as a separate library with gfx1250 filtered out). Ports the release/2.12 fix (#3458 / pytorch#190683) to release/2.13.
This was referenced Jul 27, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Enable initial gfx1250 (CDNA5) support on
release/2.13, matching whatrelease/2.12already has.release/2.13did not previously contain thegfx1250 enablement (it lacks upstream pytorch#188597 / #3421), so
this brings it in line and additionally adds the single-arch CK-GEMM guard.
Changes
Cherry-pick of [release/2.12] [ROCm] Add initial support for gfx1250 (#188597) #3421 ("[release/2.12] [ROCm] Add initial
support for gfx1250 ([ROCm] Add initial support for gfx1250 pytorch/pytorch#188597)") onto
release/2.13:ck_gemmlibrary with gfx1250 filtered out ofHIP_ARCHITECTURES(so multi-arch builds that include gfx1250 keep CK GEMMfor the other archs instead of failing to compile for gfx1250)
Single-arch CK-GEMM guard (
cmake/Dependencies.cmake): when gfx1250 isthe only arch, disable
USE_ROCM_CK_GEMMbefore the-DUSE_ROCM_CK_GEMMflag is added, so both the CK GEMM sources and their call sites compile out
consistently and single-arch gfx1250 builds link cleanly. This ports the
fix from [release/2.12] Disable CK GEMM/SDPA when gfx1250 is the only arch #3458 (release/2.12) and [ROCm] Disable CK GEMM/SDPA when gfx1250 is the only arch pytorch/pytorch#190683 (upstream main).
Companion PRs
release/2.12 + nightly. A follow-up will remove gfx1250 from the
release/2.13unsupported-family filter once this lands.Validation
TheRock dev builds for gfx1250 (single-arch + multi-arch) to be linked here.