ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-11 17:00:18 +00:00
Code Issues Packages Projects Releases Wiki Activity

[rocm-libraries] ROCm/rocm-libraries#6933 (commit ac8b7d9)

Browse Source 
[CK] Filter out unsupported targets.

## Motivation

Filter out any unsupported targets, e.g., gfx900, gfx906, gfx90c, from
the GPU_TARGETS or GPU_ARCHS lists.

## Technical Details

<!-- Explain the changes along with any relevant GitHub links. -->

## Test Plan

<!-- Explain any relevant testing done to verify this PR. -->

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
This commit is contained in:
Illia Silin
2026-05-08 03:24:33 +00:00
committed by assistant-librarian[bot]
parent 837fb379d9
commit cc29502c28
2 changed files with 32 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

19
CMakeLists.txt
View File

@@ -210,6 +210,21 @@ else()
set(USER_GPU_TARGETS 0)
endif()
#Unsupported GPU targets to be filtered from the list:
set(CK_UNSUPPORTED_GPU_TARGETS "gfx900;gfx906;gfx90c")
#If only one of the unsupported targets is requested, generate dummy target and exit here.
if("${GPU_TARGETS}" IN_LIST CK_UNSUPPORTED_GPU_TARGETS)
add_custom_target(ck_dummy_target)
message("CK is not supported for target ${GPU_TARGETS}")
return()
endif()
#If multiple targets are requested, filter out any targets currently on the unsupported list:
message(STATUS "Filtering out unsupported targets: ${CK_UNSUPPORTED_GPU_TARGETS}")
list(REMOVE_ITEM GPU_TARGETS ${CK_UNSUPPORTED_GPU_TARGETS})
list(REMOVE_ITEM GPU_ARCHS ${CK_UNSUPPORTED_GPU_TARGETS})
find_package(hip REQUIRED)
enable_language(HIP)
@@ -232,8 +247,10 @@ if(NOT ENABLE_ASAN_PACKAGING)
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1200;gfx1201")
elseif(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER_EQUAL 600400000 AND ${hip_VERSION_FLAT} LESS 600443483)
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1200;gfx1201;gfx950")
elseif(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER_EQUAL 600443483)
elseif(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER_EQUAL 600443483 AND ${hip_VERSION_FLAT} LESS 700200000)
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx950;gfx10-3-generic;gfx11-generic;gfx12-generic")
elseif(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER_EQUAL 700200000)
set(CK_GPU_TARGETS "")
endif()
else()
#build CK only for xnack-supported targets when using ASAN

20
library/src/tensor_operation_instance/gpu/CMakeLists.txt
View File

@@ -133,6 +133,14 @@ function(add_instance_library INSTANCE_NAME)
endforeach()
#set the target filters
set(UNSUPPORTED_TARGETS "gfx900;gfx906;gfx906:xnack-;gfx90c")
set(NON_XDL_TARGETS ${UNSUPPORTED_TARGETS} gfx1030 gfx1031 gfx1032 gfx1033 gfx1034 gfx1035 gfx1036 gfx10-3-generic)
set(NON_WMMA_TARGETS ${NON_XDL_TARGETS} gfx908 gfx90a gfx942 gfx950 gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack-)
set(NON_MHA_TARGETS ${NON_XDL_TARGETS} gfx908:xnack- gfx908:xnack+ gfx908 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1153 gfx1200 gfx1201 gfx11-generic gfx12-generic)
set(NON_MX_TARGETS ${NON_XDL_TARGETS} gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1153 gfx1200 gfx1201 gfx11-generic gfx12-generic)
set(NON_F8_TARGETS ${NON_XDL_TARGETS} gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1153 gfx11-generic)
message(DEBUG "remaining instances: ${ARGN}")
#only continue if there are some source files left on the list
if(ARGN)
@@ -142,25 +150,25 @@ function(add_instance_library INSTANCE_NAME)
set(INST_TARGETS ${SUPPORTED_GPU_TARGETS})
if(source_name MATCHES "_xdl")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx10-3-generic)
list(REMOVE_ITEM INST_TARGETS ${NON_XDL_TARGETS})
elseif(source_name MATCHES "_wmma")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx950)
list(REMOVE_ITEM INST_TARGETS ${NON_WMMA_TARGETS})
elseif(source_name MATCHES "mha")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx908 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1153 gfx1200 gfx1201 gfx10-3-generic gfx11-generic gfx12-generic)
list(REMOVE_ITEM INST_TARGETS ${NON_MHA_TARGETS})
endif()
if(source_name MATCHES "_mx")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1153 gfx1200 gfx1201 gfx10-3-generic gfx11-generic gfx12-generic)
list(REMOVE_ITEM INST_TARGETS ${NON_MX_TARGETS})
endif()
#only build the fp8 gemm instances for gfx90a if the build argument is set, otherwise only build for gfx942/gfx950 and gfx1200/gfx1201
if(NOT CK_USE_FP8_ON_UNSUPPORTED_ARCH)
if(source_name MATCHES "gemm_xdl_universal|gemm_multiply_multiply|gemm_universal_preshuffle|gemm_blockscale" AND source_name MATCHES "_f8_")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1153 gfx10-3-generic gfx11-generic)
list(REMOVE_ITEM INST_TARGETS ${NON_F8_TARGETS})
endif()
else()
if(source_name MATCHES "gemm_xdl_universal|gemm_multiply_multiply|gemm_universal_preshuffle|gemm_blockscale" AND source_name MATCHES "_f8_")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack- gfx908:xnack+ gfx908 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1150 gfx1151 gfx1152 gfx1153 gfx10-3-generic gfx11-generic)
list(REMOVE_ITEM INST_TARGETS ${NON_F8_TARGETS})
endif()
endif()
if(source_name MATCHES "gemm_wmma_universal" AND source_name MATCHES "f8")