From d424bbe440caa4ad4ae62af1fd877395d815f00f Mon Sep 17 00:00:00 2001 From: Illia Silin <98187287+illsilin@users.noreply.github.com> Date: Wed, 26 Mar 2025 15:15:57 -0700 Subject: [PATCH] Disable all pk_i4 tests for all targets except gfx942/950. (#2022) * only build gemm_fp8_pk_i4 examples for gfx942/950 * fix cmake logic * moved the architecture check to IsSupported function * Revert "moved the architecture check to IsSupported function" This reverts commit 056d2a08b3508cb8ffd0ebbb2b334d1fefa1cd28. * disable all pk_i4 tests for targets other than gfx942/950 * fix cmake logic [ROCm/composable_kernel commit: 23a949706cb2ab24949b98e28951c97e0f8c434e] --- Jenkinsfile | 5 +++++ example/01_gemm/CMakeLists.txt | 6 +++--- example/01_gemm/gemm_xdl_bf16_pk_i4_v3.cpp | 3 ++- example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp | 3 ++- example/01_gemm/gemm_xdl_fp16_pk_i4_v3_b_scale.cpp | 3 ++- example/65_gemm_multiply_multiply/moe_gemm1_xdl_pk_i4.cpp | 3 ++- example/65_gemm_multiply_multiply/moe_gemm2_xdl_pk_i4.cpp | 3 ++- example/CMakeLists.txt | 6 +++--- 8 files changed, 21 insertions(+), 11 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 29aec8e709..86cac3c485 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -291,6 +291,11 @@ def cmake_build(Map conf=[:]){ setup_cmd = conf.get("setup_cmd", """${cmake_envs} cmake -G Ninja ${setup_args} -DCMAKE_CXX_FLAGS=" -O3 -ftime-trace " .. """) build_cmd = conf.get("build_cmd", "${build_envs} ninja -j${nt} ${config_targets}") } + else if (setup_args.contains("gfx908;gfx90a;gfx942")){ + //limit the number of build threads when building for multiple gfx9 targets + setup_cmd = conf.get("setup_cmd", "${cmake_envs} cmake ${setup_args} .. ") + build_cmd = conf.get("build_cmd", "${build_envs} make -j32 ${config_targets}") + } else{ setup_cmd = conf.get("setup_cmd", "${cmake_envs} cmake ${setup_args} .. ") build_cmd = conf.get("build_cmd", "${build_envs} make -j${nt} ${config_targets}") diff --git a/example/01_gemm/CMakeLists.txt b/example/01_gemm/CMakeLists.txt index 09288cd4ad..ee9f959d94 100755 --- a/example/01_gemm/CMakeLists.txt +++ b/example/01_gemm/CMakeLists.txt @@ -29,9 +29,6 @@ add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_v3) add_example_executable(example_gemm_xdl_fp8_v3 gemm_xdl_fp8_v3.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp8_v3) add_example_executable(example_gemm_xdl_fp16_fp8_v3 gemm_xdl_fp16_fp8_v3.cpp) -add_example_executable(example_gemm_xdl_fp16_pk_i4_v3 gemm_xdl_fp16_pk_i4_v3.cpp) -add_example_executable(example_gemm_xdl_fp16_pk_i4_v3_b_scale gemm_xdl_fp16_pk_i4_v3_b_scale.cpp) -add_example_executable(example_gemm_xdl_bf16_pk_i4_v3 gemm_xdl_bf16_pk_i4_v3.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_fp16_fp8_v3) add_example_executable(example_gemm_xdl_bf16_v3 gemm_xdl_bf16_v3.cpp) add_example_dependencies(example_gemm_xdl example_gemm_xdl_bf16_v3) @@ -40,6 +37,9 @@ list(APPEND gpu_list gfx942 gfx950) set(target 0) foreach(gpu IN LISTS GPU_TARGETS) if(gpu IN_LIST gpu_list AND target EQUAL 0) + add_example_executable(example_gemm_xdl_fp16_pk_i4_v3 gemm_xdl_fp16_pk_i4_v3.cpp) + add_example_executable(example_gemm_xdl_fp16_pk_i4_v3_b_scale gemm_xdl_fp16_pk_i4_v3_b_scale.cpp) + add_example_executable(example_gemm_xdl_bf16_pk_i4_v3 gemm_xdl_bf16_pk_i4_v3.cpp) add_example_executable(example_gemm_xdl_fp8_pk_i4_bpreshuffle_v3 gemm_xdl_fp8_pk_i4_bpreshuffle_v3.cpp) add_example_executable(example_gemm_xdl_fp8_pk_i4_v3 gemm_xdl_fp8_pk_i4_v3.cpp) set(target 1) diff --git a/example/01_gemm/gemm_xdl_bf16_pk_i4_v3.cpp b/example/01_gemm/gemm_xdl_bf16_pk_i4_v3.cpp index 7b491173a6..7c232f1bcf 100644 --- a/example/01_gemm/gemm_xdl_bf16_pk_i4_v3.cpp +++ b/example/01_gemm/gemm_xdl_bf16_pk_i4_v3.cpp @@ -192,7 +192,8 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) b_element_op, c_element_op); - if(!gemm.IsSupportedArgument(argument)) + if(!gemm.IsSupportedArgument(argument) || ck::get_device_name() != "gfx942" || + ck::get_device_name() != "gfx950") { std::cerr << gemm.GetTypeString() << " does not support this problem" << std::endl; diff --git a/example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp b/example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp index e8a3064de6..61c5a32d5d 100644 --- a/example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp +++ b/example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp @@ -242,7 +242,8 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) b_element_op, c_element_op); - if(!gemm.IsSupportedArgument(argument)) + if(!gemm.IsSupportedArgument(argument) || ck::get_device_name() != "gfx942" || + ck::get_device_name() != "gfx950") { std::cerr << gemm.GetTypeString() << " does not support this problem" << std::endl; diff --git a/example/01_gemm/gemm_xdl_fp16_pk_i4_v3_b_scale.cpp b/example/01_gemm/gemm_xdl_fp16_pk_i4_v3_b_scale.cpp index c8a40baa8a..468dd699a1 100644 --- a/example/01_gemm/gemm_xdl_fp16_pk_i4_v3_b_scale.cpp +++ b/example/01_gemm/gemm_xdl_fp16_pk_i4_v3_b_scale.cpp @@ -274,7 +274,8 @@ bool run_gemm(const ProblemType& problem_size, const ExecutionConfig& config) b_element_op, c_element_op); - if(!gemm.IsSupportedArgument(argument)) + if(!gemm.IsSupportedArgument(argument) || ck::get_device_name() != "gfx942" || + ck::get_device_name() != "gfx950") { std::cerr << gemm.GetTypeString() << " does not support this problem" << std::endl; diff --git a/example/65_gemm_multiply_multiply/moe_gemm1_xdl_pk_i4.cpp b/example/65_gemm_multiply_multiply/moe_gemm1_xdl_pk_i4.cpp index c06e595c0f..17f4cd8a3f 100644 --- a/example/65_gemm_multiply_multiply/moe_gemm1_xdl_pk_i4.cpp +++ b/example/65_gemm_multiply_multiply/moe_gemm1_xdl_pk_i4.cpp @@ -440,7 +440,8 @@ int main(int argc, char* argv[]) b_element_op, cde_element_op); - if(!device_op.IsSupportedArgument(argument)) + if(!device_op.IsSupportedArgument(argument) || ck::get_device_name() != "gfx942" || + ck::get_device_name() != "gfx950") { throw std::runtime_error( "wrong! device_gemm with the specified compilation parameters does " diff --git a/example/65_gemm_multiply_multiply/moe_gemm2_xdl_pk_i4.cpp b/example/65_gemm_multiply_multiply/moe_gemm2_xdl_pk_i4.cpp index c80b01d8c5..8441862004 100644 --- a/example/65_gemm_multiply_multiply/moe_gemm2_xdl_pk_i4.cpp +++ b/example/65_gemm_multiply_multiply/moe_gemm2_xdl_pk_i4.cpp @@ -407,7 +407,8 @@ int main(int argc, char* argv[]) b_element_op, cde_element_op); - if(!device_op.IsSupportedArgument(argument)) + if(!device_op.IsSupportedArgument(argument) || ck::get_device_name() != "gfx942" || + ck::get_device_name() != "gfx950") { throw std::runtime_error( "wrong! device_gemm with the specified compilation parameters does " diff --git a/example/CMakeLists.txt b/example/CMakeLists.txt index 0bfba89c92..64ff2a6813 100644 --- a/example/CMakeLists.txt +++ b/example/CMakeLists.txt @@ -113,14 +113,14 @@ function(add_example_executable EXAMPLE_NAME FILE_NAME) endforeach() #only continue if there are some source files left on the list if(FILE_NAME) - if(FILE_NAME MATCHES "_xdl" AND NOT FILE_NAME MATCHES "_fp8_pk_i4") + if(FILE_NAME MATCHES "_xdl" AND NOT FILE_NAME MATCHES "_pk_i4") list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic) elseif(FILE_NAME MATCHES "_wmma") list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx950) elseif(FILE_NAME MATCHES "_mx") #only build mx example for gfx950 list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx942 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic) - elseif(FILE_NAME MATCHES "_fp8_pk_i4") #only build these examples for gfx942 and gfx950 - message("trimming targets for ${FILE_NAME}") + elseif(FILE_NAME MATCHES "_pk_i4") #only build these examples for gfx942 and gfx950 + message("trimming targets for ${FILE_NAME}") list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx906:xnack- gfx908:xnack+ gfx908:xnack- gfx90a:xnack+ gfx90a:xnack- gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201 gfx10.3-generic gfx11-generic gfx12-generic) endif() set_source_files_properties(${FILE_NAME} PROPERTIES LANGUAGE HIP)