Skip to content

[release/2.12] [ROCm] Add initial support for gfx1250 (#188597)#3421

Merged
pragupta merged 2 commits into
release/2.12from
release_2.12_gfx1250
Jul 9, 2026
Merged

[release/2.12] [ROCm] Add initial support for gfx1250 (#188597)#3421
pragupta merged 2 commits into
release/2.12from
release_2.12_gfx1250

Conversation

@rraminen

@rraminen rraminen commented Jul 9, 2026

Copy link
Copy Markdown

Summary

Validation theRock Run: https://github.com/ROCm/TheRock/actions/runs/29045360944

Cherry-pick of upstream pytorch/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:

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

rraminen and others added 2 commits July 9, 2026 18:15
This PR adds initial ROCm support for gfx1250 (CDNA5).

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

Pull Request resolved: pytorch#188597
Approved by: https://github.com/jeffdaily, https://github.com/pruthvistony

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>
@rocm-repo-management-api

Copy link
Copy Markdown

Jenkins build for 6722f8ea68735ee20455fffae1907038ff0185f9 commit is in progress
Links: Pipeline Overview / Build artifacts

@jithunnair-amd jithunnair-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.

LGTM, but please verify build and test changes via theRock CI to ensure we are not introducing breakages.

@pragupta pragupta merged commit 13eab2b into release/2.12 Jul 9, 2026
0 of 2 checks passed
@pragupta pragupta deleted the release_2.12_gfx1250 branch July 9, 2026 21:39
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants