From b17fa5656f2d5c04586b5962eadc49359c2b5e0a Mon Sep 17 00:00:00 2001 From: Thrupti Raj Lakshmana Gowda Date: Wed, 24 Dec 2025 12:45:56 -0600 Subject: [PATCH] [CK TILE ENGINE] CI configuration with basic cases (#3475) * [CK TILE ENGINE] Adding GEMM BASIC TEST in Kenkins * fix RUN_TILE_ENGINE_BASIC_TESTS name typo * [CK Tile Engine] Updating basic CI * Resolving merging issues * Resolving merging issues --------- Co-authored-by: illsilin_amdeng [ROCm/composable_kernel commit: 62a8ec155facd901232977b688d5225d72969709] --- Jenkinsfile | 45 +++++++++- .../configs/default_ci_config.json | 89 +++++++++++++++++++ .../configs/default_ci_config.json | 89 +++++++++++++++++++ .../configs/default_ci_config.json | 89 +++++++++++++++++++ 4 files changed, 311 insertions(+), 1 deletion(-) create mode 100644 tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config.json create mode 100644 tile_engine/ops/gemm/gemm_preshuffle/configs/default_ci_config.json create mode 100644 tile_engine/ops/gemm/gemm_universal/configs/default_ci_config.json diff --git a/Jenkinsfile b/Jenkinsfile index f90eedf107..cb2f8631c5 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -1114,7 +1114,7 @@ def run_pytorch_tests(Map conf=[:]){ //launch develop branch daily jobs CRON_SETTINGS = BRANCH_NAME == "develop" ? '''0 23 * * * % RUN_FULL_QA=true;RUN_CK_TILE_FMHA_TESTS=true;RUN_PERFORMANCE_TESTS=true;FORCE_CI=true - 0 22 * * * % RUN_FULL_QA=true;DISABLE_DL_KERNELS=true;RUN_TILE_ENGINE_GEMM_TESTS=true;RUN_PERFORMANCE_TESTS=true;RUN_ALL_UNIT_TESTS=true;FORCE_CI=true + 0 22 * * * % RUN_FULL_QA=true;DISABLE_DL_KERNELS=true;RUN_TILE_ENGINE_BASIC_TESTS=true;RUN_TILE_ENGINE_GEMM_TESTS=true;RUN_PERFORMANCE_TESTS=true;RUN_ALL_UNIT_TESTS=true;FORCE_CI=true 0 21 * * * % RUN_GROUPED_CONV_LARGE_CASES_TESTS=true;hipTensor_test=true;BUILD_GFX101=false;BUILD_GFX908=false;BUILD_GFX942=true;BUILD_GFX950=true;RUN_PERFORMANCE_TESTS=true;RUN_ALL_UNIT_TESTS=true;FORCE_CI=true;BUILD_PACKAGES=true 0 19 * * * % BUILD_DOCKER=true;COMPILER_VERSION=amd-staging;BUILD_COMPILER=/llvm-project/build/bin/clang++;USE_SCCACHE=false;NINJA_BUILD_TRACE=true;RUN_ALL_UNIT_TESTS=true;FORCE_CI=true 0 17 * * * % BUILD_DOCKER=true;COMPILER_VERSION=amd-mainline;BUILD_COMPILER=/llvm-project/build/bin/clang++;USE_SCCACHE=false;NINJA_BUILD_TRACE=true;RUN_ALL_UNIT_TESTS=true;FORCE_CI=true @@ -1199,6 +1199,10 @@ pipeline { name: "RUN_CK_TILE_FMHA_TESTS", defaultValue: false, description: "Run the ck_tile FMHA tests (default: OFF)") + booleanParam( + name: "RUN_TILE_ENGINE_BASIC_TESTS", + defaultValue: false, + description: "Run the tile_engine_basic tests (default: OFF)") booleanParam( name: "RUN_TILE_ENGINE_GEMM_TESTS", defaultValue: false, @@ -1616,6 +1620,45 @@ pipeline { } } } + stage("Run TILE_ENGINE_BASIC Tests") + { + when { + beforeAgent true + expression { env.SHOULD_RUN_CI.toBoolean() } + } + parallel + { + stage("Run TILE_ENGINE_BASIC Tests on gfx942") + { + when { + beforeAgent true + expression { params.RUN_TILE_ENGINE_BASIC_TESTS.toBoolean() } + } + agent{ label rocmnode("gfx942") } + environment{ + setup_args = "NO_CK_BUILD" + execute_args = """ cmake -G Ninja -D CMAKE_PREFIX_PATH=/opt/rocm \ + -D CMAKE_CXX_COMPILER="${params.BUILD_COMPILER}" \ + -D CMAKE_BUILD_TYPE=Release \ + -D GPU_TARGETS="gfx942" \ + -D GEMM_UNIVERSAL_DATATYPE="fp8;fp16" \ + -D GEMM_UNIVERSAL_LAYOUT="rcr;rrr;crr;ccr" \ + -D GEMM_UNIVERSAL_CONFIG_FILE="default_ci_config.json" \ + -D GEMM_MULTI_D_DATATYPE="fp16" \ + -D GEMM_MULTI_D_LAYOUT="rcrr;rrrr;crrr;ccrr" \ + -D GEMM_MULTI_D_CONFIG_FILE="default_ci_config.json" \ + -D GEMM_PRESHUFFLE_DATATYPE="fp16;fp8;bf16;bf8" \ + -D GEMM_PRESHUFFLE_LAYOUT="rcr" \ + -D GEMM_PRESHUFFLE_CONFIG_FILE="default_ci_config.json" .. && \ + ninja -j${nthreads()} benchmark_gemm_universal_all benchmark_gemm_preshuffle_all benchmark_gemm_multi_d_all """ + } + steps{ + buildHipClangJobAndReboot(setup_args:setup_args, build_type: 'Release', execute_cmd: execute_args) + cleanWs() + } + } + } + } stage("Run TILE_ENGINE_GEMM Tests") { when { diff --git a/tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config.json b/tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config.json new file mode 100644 index 0000000000..2df04d0ac1 --- /dev/null +++ b/tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config.json @@ -0,0 +1,89 @@ +{ + "tile_config": { + "tile_m": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_n": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_k": { + "max": 256, + "min": 64, + "step": 256 + }, + "warp_m": { + "values": [ + 2 + ] + }, + "warp_n": { + "values": [ + 2 + ] + }, + "warp_k": { + "values": [ + 1 + ] + }, + "warp_tile_m": { + "values": [ + 32 + ] + }, + "warp_tile_n": { + "values": [ + 32 + ] + }, + "warp_tile_k": { + "values": [ + 64 + ] + } + }, + "trait_config": { + "pipeline": { + "values": [ + "compv4" + ] + }, + "scheduler": { + "values": [ + "intrawave", + "interwave" + ] + }, + "epilogue": { + "values": [ + "cshuffle" + ] + }, + "pad_m": { + "values": [ + false + ] + }, + "pad_n": { + "values": [ + false + ] + }, + "pad_k": { + "values": [ + false + ] + }, + "persistent": { + "values": [ + false, + true + ] + } + }, + "k_block_per_cu": 1 +} diff --git a/tile_engine/ops/gemm/gemm_preshuffle/configs/default_ci_config.json b/tile_engine/ops/gemm/gemm_preshuffle/configs/default_ci_config.json new file mode 100644 index 0000000000..b32d42dabf --- /dev/null +++ b/tile_engine/ops/gemm/gemm_preshuffle/configs/default_ci_config.json @@ -0,0 +1,89 @@ +{ + "tile_config": { + "tile_m": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_n": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_k": { + "max": 256, + "min": 64, + "step": 256 + }, + "warp_m": { + "values": [ + 2 + ] + }, + "warp_n": { + "values": [ + 2 + ] + }, + "warp_k": { + "values": [ + 1 + ] + }, + "warp_tile_m": { + "values": [ + 16 + ] + }, + "warp_tile_n": { + "values": [ + 16 + ] + }, + "warp_tile_k": { + "values": [ + 32 + ] + } + }, + "trait_config": { + "pipeline": { + "values": [ + "preshufflev2" + ] + }, + "scheduler": { + "values": [ + "default" + ] + }, + "epilogue": { + "values": [ + "cshuffle" + ] + }, + "pad_m": { + "values": [ + false + ] + }, + "pad_n": { + "values": [ + false + ] + }, + "pad_k": { + "values": [ + false + ] + }, + "persistent": { + "values": [ + true, + false + ] + } + }, + "k_block_per_cu": 1, + "permute_n": true +} \ No newline at end of file diff --git a/tile_engine/ops/gemm/gemm_universal/configs/default_ci_config.json b/tile_engine/ops/gemm/gemm_universal/configs/default_ci_config.json new file mode 100644 index 0000000000..2dd8230edc --- /dev/null +++ b/tile_engine/ops/gemm/gemm_universal/configs/default_ci_config.json @@ -0,0 +1,89 @@ +{ + "tile_config": { + "tile_m": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_n": { + "max": 256, + "min": 64, + "step": 256 + }, + "tile_k": { + "max": 256, + "min": 64, + "step": 256 + }, + "warp_m": { + "values": [ + 2 + ] + }, + "warp_n": { + "values": [ + 2 + ] + }, + "warp_k": { + "values": [ + 1 + ] + }, + "warp_tile_m": { + "values": [ + 16 + ] + }, + "warp_tile_n": { + "values": [ + 16 + ] + }, + "warp_tile_k": { + "values": [ + 32 + ] + } + }, + "trait_config": { + "pipeline": { + "values": [ + "compv4" + ] + }, + "scheduler": { + "values": [ + "intrawave", + "interwave" + ] + }, + "epilogue": { + "values": [ + "cshuffle" + ] + }, + "pad_m": { + "values": [ + false + ] + }, + "pad_n": { + "values": [ + false + ] + }, + "pad_k": { + "values": [ + false + ] + }, + "persistent": { + "values": [ + false, + true + ] + } + }, + "k_block_per_cu": 1 +}