From 0a5c8d6e450ddcb9c875cbdae0b4a3706f26fab3 Mon Sep 17 00:00:00 2001 From: Vidyasagar Ananthan Date: Tue, 19 Aug 2025 18:12:06 -0700 Subject: [PATCH] Setting gpu target filtering for tile engine to gfx90a, gfx942 and gfx950. (#2709) [ROCm/composable_kernel commit: bf3e719c16846c704e8b93b0116954b321933d74] --- tile_engine/ops/gemm/CMakeLists.txt | 21 +++++++++++++++++++++ tile_engine/ops/gemm_multi_d/CMakeLists.txt | 21 +++++++++++++++++++++ 2 files changed, 42 insertions(+) diff --git a/tile_engine/ops/gemm/CMakeLists.txt b/tile_engine/ops/gemm/CMakeLists.txt index fe9b7802a7..42c114b499 100644 --- a/tile_engine/ops/gemm/CMakeLists.txt +++ b/tile_engine/ops/gemm/CMakeLists.txt @@ -3,6 +3,24 @@ set(GEMM_DATATYPE "fp8;fp16" CACHE STRING "List of datatypes for GEMM (semicolon set(GEMM_LAYOUT "rcr" CACHE STRING "List of layout for GEMM (semicolon-separated)") function(build_gemm_for_datatype datatype layout) + # Filter GPU targets to only gfx90a, gfx942, and gfx950 + set(GEMM_GPU_TARGETS "") + set(DESIRED_TARGETS "gfx90a;gfx942;gfx950") + + foreach(target IN LISTS SUPPORTED_GPU_TARGETS) + if(target IN_LIST DESIRED_TARGETS) + list(APPEND GEMM_GPU_TARGETS ${target}) + endif() + endforeach() + + # Skip compilation if no matching targets found + if(NOT GEMM_GPU_TARGETS) + message(WARNING "Skipping Tile Engine GEMM compilation: No supported GPU targets (gfx90a, gfx942, gfx950) found in SUPPORTED_GPU_TARGETS: ${SUPPORTED_GPU_TARGETS}") + return() + endif() + + message(STATUS "Building GEMM for GPU targets: ${GEMM_GPU_TARGETS}") + set(working_path "${CMAKE_CURRENT_BINARY_DIR}/${datatype}/${layout}") # Comment this if-else block when using user_provided_config @@ -83,6 +101,7 @@ function(build_gemm_for_datatype datatype layout) if(chunk_files) set(sub_intermediate_lib_name "gemm_objlib_${name}_${i}_${datatype}_${layout}") add_library(${sub_intermediate_lib_name} OBJECT ${chunk_files}) + set_property(TARGET ${sub_intermediate_lib_name} PROPERTY HIP_ARCHITECTURES ${GEMM_GPU_TARGETS}) list(APPEND sub_intermediate_libs ${sub_intermediate_lib_name}) endif() @@ -102,6 +121,7 @@ function(build_gemm_for_datatype datatype layout) add_library(${intermediate_lib_name} STATIC ${obj_exprs}) add_dependencies(${intermediate_lib_name} gemm_gen_${datatype}_${layout}) + set_property(TARGET ${intermediate_lib_name} PROPERTY HIP_ARCHITECTURES ${GEMM_GPU_TARGETS}) #foreach(objlib IN LISTS sub_intermediate_libs) # target_sources(${intermediate_lib_name} PRIVATE $) #endforeach() @@ -132,6 +152,7 @@ function(build_gemm_for_datatype datatype layout) # Executable per datatype set(exec_name "benchmark_gemm_${datatype}_${layout}") add_executable(${exec_name} benchmark_gemm.cpp) + set_property(TARGET ${exec_name} PROPERTY HIP_ARCHITECTURES ${GEMM_GPU_TARGETS}) target_link_libraries(${exec_name} PRIVATE gemm_host_api_${datatype}_${layout}) target_compile_options(${exec_name} PRIVATE -Wno-undefined-func-template diff --git a/tile_engine/ops/gemm_multi_d/CMakeLists.txt b/tile_engine/ops/gemm_multi_d/CMakeLists.txt index 3708dd3fee..dc08e9cad3 100644 --- a/tile_engine/ops/gemm_multi_d/CMakeLists.txt +++ b/tile_engine/ops/gemm_multi_d/CMakeLists.txt @@ -4,6 +4,24 @@ set(GEMM_MULTI_D_LAYOUT "rcrr" CACHE STRING "List of layout for GEMM Multi D(sem set(GEMM_MULTI_D_ELEMENTWISE_FUNCTION "mul" CACHE STRING "Elementwise function") function(build_gemm_multi_d_for_datatype_layout datatype layout) + # Filter GPU targets to only gfx90a, gfx942, and gfx950 + set(GEMM_GPU_TARGETS "") + set(DESIRED_TARGETS "gfx90a;gfx942;gfx950") + + foreach(target IN LISTS SUPPORTED_GPU_TARGETS) + if(target IN_LIST DESIRED_TARGETS) + list(APPEND GEMM_GPU_TARGETS ${target}) + endif() + endforeach() + + # Skip compilation if no matching targets found + if(NOT GEMM_GPU_TARGETS) + message(WARNING "Skipping Tile Engine GEMM Multi D compilation: No supported GPU targets (gfx90a, gfx942, gfx950) found in SUPPORTED_GPU_TARGETS: ${SUPPORTED_GPU_TARGETS}") + return() + endif() + + message(STATUS "Building GEMM Multi D for GPU targets: ${GEMM_GPU_TARGETS}") + set(working_path "${CMAKE_CURRENT_BINARY_DIR}/${datatype}/${layout}") # Comment this if-else block when using user_provided_config @@ -86,6 +104,7 @@ function(build_gemm_multi_d_for_datatype_layout datatype layout) if(chunk_files) set(sub_intermediate_lib_name "gemm_multi_d_objlib_${name}_${i}_${datatype}_${layout}") add_library(${sub_intermediate_lib_name} OBJECT ${chunk_files}) + set_property(TARGET ${sub_intermediate_lib_name} PROPERTY HIP_ARCHITECTURES ${GEMM_GPU_TARGETS}) list(APPEND sub_intermediate_libs ${sub_intermediate_lib_name}) endif() @@ -105,6 +124,7 @@ function(build_gemm_multi_d_for_datatype_layout datatype layout) add_library(${intermediate_lib_name} STATIC ${obj_exprs}) add_dependencies(${intermediate_lib_name} gemm_multi_d_gen_${datatype}_${layout}) + set_property(TARGET ${intermediate_lib_name} PROPERTY HIP_ARCHITECTURES ${GEMM_GPU_TARGETS}) #foreach(objlib IN LISTS sub_intermediate_libs) # target_sources(${intermediate_lib_name} PRIVATE $) #endforeach() @@ -136,6 +156,7 @@ function(build_gemm_multi_d_for_datatype_layout datatype layout) # Executable per datatype set(exec_name "benchmark_gemm_multi_d_${datatype}_${layout}") add_executable(${exec_name} benchmark_gemm_multi_d.cpp) + set_property(TARGET ${exec_name} PROPERTY HIP_ARCHITECTURES ${GEMM_GPU_TARGETS}) target_link_libraries(${exec_name} PRIVATE gemm_multi_d_host_api_${datatype}_${layout}) target_compile_options(${exec_name} PRIVATE -Wno-undefined-func-template