From 3731b213435b34ccbd9d1d8950426832087ec5aa Mon Sep 17 00:00:00 2001 From: ethanwee1 Date: Tue, 21 Jul 2026 16:36:31 +0000 Subject: [PATCH 1/2] [ROCm] Disable CK GEMM when gfx1250 is the only arch 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. --- aten/src/ATen/CMakeLists.txt | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/aten/src/ATen/CMakeLists.txt b/aten/src/ATen/CMakeLists.txt index fc5d69381e648..9f641907b88b0 100644 --- a/aten/src/ATen/CMakeLists.txt +++ b/aten/src/ATen/CMakeLists.txt @@ -684,6 +684,12 @@ if(USE_ROCM) ) file(GLOB native_hip_bgemm CONFIGURE_DEPENDS "native/hip/bgemm_kernels/*.hip") file(GLOB native_hip_ck CONFIGURE_DEPENDS "native/hip/ck*.hip") + # composable_kernel has no gfx1250 support yet. If gfx1250 is the only arch, + # disable CK GEMM so we do not create ck_gemm with empty HIP_ARCHITECTURES. + 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 (composable_kernel lacks gfx1250 support)") + caffe2_update_option(USE_ROCM_CK_GEMM OFF) + endif() if(NOT USE_ROCM_CK_GEMM) exclude(ATen_HIP_SRCS "${ATen_HIP_SRCS}" ${native_hip_bgemm} ${native_hip_ck}) From e7167bb379238dbf63400713cc2034a6e92c778c Mon Sep 17 00:00:00 2001 From: ethanwee1 Date: Wed, 22 Jul 2026 15:55:17 +0000 Subject: [PATCH 2/2] [ROCm] Move gfx1250 CK GEMM disable before -DUSE_ROCM_CK_GEMM flag 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. --- aten/src/ATen/CMakeLists.txt | 6 ------ cmake/Dependencies.cmake | 10 ++++++++++ 2 files changed, 10 insertions(+), 6 deletions(-) diff --git a/aten/src/ATen/CMakeLists.txt b/aten/src/ATen/CMakeLists.txt index 9f641907b88b0..fc5d69381e648 100644 --- a/aten/src/ATen/CMakeLists.txt +++ b/aten/src/ATen/CMakeLists.txt @@ -684,12 +684,6 @@ if(USE_ROCM) ) file(GLOB native_hip_bgemm CONFIGURE_DEPENDS "native/hip/bgemm_kernels/*.hip") file(GLOB native_hip_ck CONFIGURE_DEPENDS "native/hip/ck*.hip") - # composable_kernel has no gfx1250 support yet. If gfx1250 is the only arch, - # disable CK GEMM so we do not create ck_gemm with empty HIP_ARCHITECTURES. - 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 (composable_kernel lacks gfx1250 support)") - caffe2_update_option(USE_ROCM_CK_GEMM OFF) - endif() if(NOT USE_ROCM_CK_GEMM) exclude(ATen_HIP_SRCS "${ATen_HIP_SRCS}" ${native_hip_bgemm} ${native_hip_ck}) diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index 70c5d3de4cf23..0c861ed515f26 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -1025,6 +1025,16 @@ if(USE_ROCM) if(HIPBLASLT_VEC_EXT) list(APPEND HIP_CXX_FLAGS -DHIPBLASLT_VEC_EXT) endif() + # composable_kernel has no gfx1250 support yet. If gfx1250 is the only arch, + # disable CK GEMM here (before -DUSE_ROCM_CK_GEMM is added) so that both the + # CK GEMM sources and their call sites (guarded by USE_ROCM_CK_GEMM in + # CUDABlas.cpp/GroupedBlas.cpp) are compiled out consistently. Doing this in + # aten/src/ATen/CMakeLists.txt instead excludes the sources but still defines + # USE_ROCM_CK_GEMM, leaving undefined references to gemm_internal_ck at link. + if("${PYTORCH_ROCM_ARCH}" STREQUAL "gfx1250") + message(WARNING "gfx1250 is the only arch in PYTORCH_ROCM_ARCH: disabling USE_ROCM_CK_GEMM (composable_kernel lacks gfx1250 support)") + caffe2_update_option(USE_ROCM_CK_GEMM OFF) + endif() if(USE_ROCM_CK_GEMM) list(APPEND HIP_CXX_FLAGS -DUSE_ROCM_CK_GEMM) endif()