Skip to content

[release/2.12] Disable CK GEMM/SDPA when gfx1250 is the only arch#3458

Draft
ethanwee1 wants to merge 2 commits into
ROCm:release/2.12from
ethanwee1:ethanwee/fix-gfx1250-ck-gemm-rocm-2.12
Draft

[release/2.12] Disable CK GEMM/SDPA when gfx1250 is the only arch#3458
ethanwee1 wants to merge 2 commits into
ROCm:release/2.12from
ethanwee1:ethanwee/fix-gfx1250-ck-gemm-rocm-2.12

Conversation

@ethanwee1

@ethanwee1 ethanwee1 commented Jul 21, 2026

Copy link
Copy Markdown

No description provided.

@ethanwee1 ethanwee1 mentioned this pull request Jul 21, 2026
Merged
@rocm-repo-management-api

rocm-repo-management-api Bot commented Jul 21, 2026 •
edited
Loading

Copy link
Copy Markdown

Jenkins build for 0382a55233af415242de97c69644905524ea6536 commit finished as ABORTED
Links: Pipeline Overview / Build artifacts / Test Results

jithunnair-amd
jithunnair-amd reviewed Jul 22, 2026
View reviewed changes
Comment thread aten/src/ATen/CMakeLists.txt Outdated
Comment on lines +208 to +216
if(USE_ROCM AND "gfx1250" IN_LIST PYTORCH_ROCM_ARCH)
set(_ck_supported_archs ${PYTORCH_ROCM_ARCH})
list(REMOVE_ITEM _ck_supported_archs gfx1250)
if("${_ck_supported_archs}" STREQUAL "")
message(WARNING "gfx1250 is the only arch in PYTORCH_ROCM_ARCH: disabling USE_ROCM_CK_GEMM and USE_ROCM_CK_SDPA (composable_kernel lacks gfx1250 support)")
caffe2_update_option(USE_ROCM_CK_GEMM OFF)
caffe2_update_option(USE_ROCM_CK_SDPA OFF)
endif()
endif()

@jithunnair-amd jithunnair-amd Jul 22, 2026

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.

Suggested change
if(USE_ROCM AND "gfx1250" IN_LIST PYTORCH_ROCM_ARCH)
set(_ck_supported_archs ${PYTORCH_ROCM_ARCH})
list(REMOVE_ITEM _ck_supported_archs gfx1250)
if("${_ck_supported_archs}" STREQUAL "")
message(WARNING "gfx1250 is the only arch in PYTORCH_ROCM_ARCH: disabling USE_ROCM_CK_GEMM and USE_ROCM_CK_SDPA (composable_kernel lacks gfx1250 support)")
caffe2_update_option(USE_ROCM_CK_GEMM OFF)
caffe2_update_option(USE_ROCM_CK_SDPA OFF)
endif()
endif()
if(USE_ROCM AND "${PYTORCH_ROCM_ARCH}" STREQUAL "gfx1250")
message(WARNING "gfx1250 is the only arch in PYTORCH_ROCM_ARCH: disabling USE_ROCM_CK_GEMM and USE_ROCM_CK_SDPA (composable_kernel lacks gfx1250 support)")
caffe2_update_option(USE_ROCM_CK_GEMM OFF)
caffe2_update_option(USE_ROCM_CK_SDPA OFF)
endif()

@jithunnair-amd jithunnair-amd Jul 22, 2026

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.

@ethanwee1 Apart from the simplification above, I'd suggest we have this logic reside next to https://github.com/ROCm/pytorch/pull/3458/changes#diff-ce80f3115ab2f6be5142f0678a1fc92c6b2d7727766ce44f48726c99e720f777R713, where we already have the arch filtering logic for gfx1250. That should update the option to OFF and skip all the subsequent code to add CK GEMM sources if gfx1250 is the only target.

The USE_ROCM_CK_SDPA logic is already set up to update the option to OFF if no supported arch is in PYTORCH_ROCM_ARCH.

ethanwee1 added a commit to ethanwee1/rocm-pytorch-ethanwee that referenced this pull request Jul 22, 2026
@ethanwee1
[ROCm] Disable CK GEMM when gfx1250 is the only arch
c8443b5 
If PYTORCH_ROCM_ARCH is gfx1250-only, composable_kernel has no supported
arch for ck_gemm and CMake fails with empty HIP_ARCHITECTURES. Disable
USE_ROCM_CK_GEMM next to the existing gfx1250 ck_gemm filtering logic.

Per review on ROCm#3458.
@ethanwee1
[ROCm] Disable CK GEMM when gfx1250 is the only arch
3731b21 
Move the gfx1250-only guard next to the existing ck_gemm arch filtering and
simplify to a single-arch check. Per review by @jithunnair-amd.
@ethanwee1

ethanwee1 commented Jul 22, 2026

Copy link
Copy Markdown
Author

Addressed review feedback from @jithunnair-amd:

  • Removed the early guard block near flash-attention setup
  • Added simplified gfx1250-only check next to the existing `ck_gemm` arch filtering (`"${PYTORCH_ROCM_ARCH}" STREQUAL "gfx1250"`)
  • Only disables `USE_ROCM_CK_GEMM` here (SDPA already auto-disables when no gfx942/gfx950)

Same change applied to upstream companion pytorch#190683

@ethanwee1
ethanwee1 force-pushed the ethanwee/fix-gfx1250-ck-gemm-rocm-2.12 branch from 0382a55 to 3731b21 Compare July 22, 2026 15:34
@ethanwee1 ethanwee1 mentioned this pull request Jul 22, 2026
Draft
@rocm-repo-management-api

rocm-repo-management-api Bot commented Jul 22, 2026 •
edited
Loading

Copy link
Copy Markdown

Jenkins build for e7167bb379238dbf63400713cc2034a6e92c778c commit finished as NOT_BUILT
Links: Pipeline Overview / Build artifacts / Test Results

@ethanwee1
[ROCm] Move gfx1250 CK GEMM disable before -DUSE_ROCM_CK_GEMM flag
e7167bb 
Disabling USE_ROCM_CK_GEMM in aten/src/ATen/CMakeLists.txt excluded the CK
GEMM source files but still left -DUSE_ROCM_CK_GEMM defined (it is emitted
earlier in cmake/Dependencies.cmake). The call sites in CUDABlas.cpp and
GroupedBlas.cpp are guarded by USE_ROCM_CK_GEMM, so they were still compiled
and linked against gemm_internal_ck/bgemm_internal_ck/group_gemm_ck, producing
undefined references at link time for gfx1250-only builds.

Move the gfx1250-only disable into cmake/Dependencies.cmake, before the
-DUSE_ROCM_CK_GEMM flag is added, so the sources and their call sites are
compiled out consistently.
@rocm-repo-management-api

rocm-repo-management-api Bot commented Jul 22, 2026 •
edited
Loading

Copy link
Copy Markdown

Jenkins build for e7167bb379238dbf63400713cc2034a6e92c778c commit finished as ABORTED
Links: Pipeline Overview / Build artifacts / Test Results

ethanwee1 added a commit to ROCm/TheRock that referenced this pull request Jul 22, 2026
@ethanwee1
build(pytorch): enable gfx1250 wheels for release/2.12 and nightly (#…
0f2827e 
…6735)

JIRA ID : AIPYTORCH-804

## Summary

Enable `gfx125X-dcgpu` (gfx1250) in TheRock's PyTorch release matrix for
**`release/2.12`** and **`nightly`**, matching the existing
`release/2.11` behavior.

Previously, `configure_pytorch_release_matrix.py` filtered gfx1250 out
of those refs even though PyTorch now supports the arch:

- `release/2.12`: ROCm/pytorch#3421
- `nightly`: upstream pytorch/pytorch#188597 (via TheRock's upstream
nightly checkout)

## Changes

- Remove `gfx125X-dcgpu` from the unsupported-family filter for
`release/2.12` and `nightly` in `configure_pytorch_release_matrix.py`.

## Validation

Both the single-arch and multi-arch gfx1250 PyTorch wheel builds are
green with this matrix change (Python 3.12, ROCm `7.15.0a20260712`,
index `rocm.nightlies.amd.com/whl-multi-arch`,
`test_amdgpu_families=none`):

| Scenario | `amdgpu_families` | Run | Result |
|----------|-------------------|-----|--------|
| Single-arch (gfx1250 only) | `gfx125X-dcgpu` |
https://github.com/ROCm/TheRock/actions/runs/29936052631 | ✅ success |
| Multi-arch (gfx942 + gfx1250) | `gfx94X-dcgpu;gfx125X-dcgpu` |
https://github.com/ROCm/TheRock/actions/runs/29938470834 | ✅ success |

Notes:

- The **multi-arch** build (TheRock's production release/nightly path)
works on stock `release/2.12` and `main` because `ck_gemm` is built as a
separate library with gfx1250 filtered out of `HIP_ARCHITECTURES`;
gfx1250 falls back to hipBLASLt at runtime, and CK GEMM stays enabled
for gfx942/gfx950.
- The **single-arch** gfx1250-only case additionally requires the
CK-GEMM CMake fix in ROCm/pytorch#3458 (`release/2.12`) and
pytorch/pytorch#190683 (`main`/nightly). These are not required for the
multi-arch production builds.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@jithunnair-amd jithunnair-amd jithunnair-amd left review comments

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@ethanwee1 @jithunnair-amd