From 17e7a733f953d79b9632debd0cf24d493cb8f6cf Mon Sep 17 00:00:00 2001 From: ThruptiRajLakshmanaGowda Date: Tue, 13 Jan 2026 03:28:37 +0000 Subject: [PATCH 1/6] memory op changes --- tile_engine/ops/gemm/gemm_instance_builder.py | 55 +++++++------------ 1 file changed, 20 insertions(+), 35 deletions(-) diff --git a/tile_engine/ops/gemm/gemm_instance_builder.py b/tile_engine/ops/gemm/gemm_instance_builder.py index 089f9686490..3046f6658fc 100644 --- a/tile_engine/ops/gemm/gemm_instance_builder.py +++ b/tile_engine/ops/gemm/gemm_instance_builder.py @@ -643,40 +643,40 @@ def populate_launch( using GemmPipeline = {pipeline_impl_map.get(pipeline)};""" - # Runfunction body - instance_code += """ + # # Runfunction body + # instance_code += """ - const auto Run = [&](const auto memory_operation_) {""" + # const auto Run = [&](const auto memory_operation_) {""" # Scheduler initialization if self.kernel_name_prefix in ["gemm_universal"]: instance_code += f""" - constexpr auto scheduler = {scheduler_type_map.get(scheduler)};""" + constexpr auto scheduler = {scheduler_type_map.get(scheduler)};""" - # Memory operation - instance_code += """ - [[maybe_unused]] constexpr auto memory_operation = memory_operation_.value;""" + # # Memory operation + # instance_code += """ + # [[maybe_unused]] constexpr auto memory_operation = memory_operation_.value;""" # UniversalGemmProblem if self.kernel_name_prefix in ["gemm_universal"]: instance_code += """ - using UniversalGemmProblem = ck_tile::UniversalGemmPipelineProblem< - ADataType, - BDataType, - AccDataType, - TileShape, - ck_tile::TileGemmUniversalTraits, - scheduler>;""" + using UniversalGemmProblem = ck_tile::UniversalGemmPipelineProblem< + ADataType, + BDataType, + AccDataType, + TileShape, + ck_tile::TileGemmUniversalTraits, + scheduler>;""" # GemmPipeline if self.kernel_name_prefix in ["gemm_universal"]: instance_code += f""" - using GemmPipeline = {pipeline_impl_map.get(pipeline)};""" + using GemmPipeline = {pipeline_impl_map.get(pipeline)};""" # Epilogue instance_code += self.populate_epilogue(epilogue) @@ -748,23 +748,8 @@ def populate_launch( ck_tile::make_kernel(GemmKernel{{}}, grids, blocks, 0, kargs)); return ave_time; - }};""" - - # Run SplitK handler - - instance_code += """ - - float ave_time = 0.f; - if(args.k_batch == 1) { - ave_time = Run(ck_tile::integral_constant{}); - } else { - ave_time = Run(ck_tile::integral_constant{}); - } - return ave_time; - } -}; + }} +}}; """ return instance_code From b7c960cd3df758d9b7488057f5dbe477f5d607b4 Mon Sep 17 00:00:00 2001 From: ThruptiRajLakshmanaGowda Date: Tue, 13 Jan 2026 04:03:44 +0000 Subject: [PATCH 2/6] memory op changes --- tile_engine/ops/gemm/gemm_instance_builder.py | 9 --------- 1 file changed, 9 deletions(-) diff --git a/tile_engine/ops/gemm/gemm_instance_builder.py b/tile_engine/ops/gemm/gemm_instance_builder.py index 3046f6658fc..9c60c565de5 100644 --- a/tile_engine/ops/gemm/gemm_instance_builder.py +++ b/tile_engine/ops/gemm/gemm_instance_builder.py @@ -643,20 +643,11 @@ def populate_launch( using GemmPipeline = {pipeline_impl_map.get(pipeline)};""" - # # Runfunction body - # instance_code += """ - - # const auto Run = [&](const auto memory_operation_) {""" - # Scheduler initialization if self.kernel_name_prefix in ["gemm_universal"]: instance_code += f""" constexpr auto scheduler = {scheduler_type_map.get(scheduler)};""" - # # Memory operation - # instance_code += """ - # [[maybe_unused]] constexpr auto memory_operation = memory_operation_.value;""" - # UniversalGemmProblem if self.kernel_name_prefix in ["gemm_universal"]: instance_code += """ From 17a02a3e3cae28a3a2b2f1cf3954e178a60fa251 Mon Sep 17 00:00:00 2001 From: ThruptiRajLakshmanaGowda Date: Tue, 13 Jan 2026 17:58:51 +0000 Subject: [PATCH 3/6] Fixing TILE_ENGINE_BASIC in Tile Engine --- Jenkinsfile | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/Jenkinsfile b/Jenkinsfile index 7292d9b70ca..91c537d8b92 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -1650,7 +1650,10 @@ pipeline { -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 """ + ninja -j${nthreads()} benchmark_gemm_universal_all benchmark_gemm_preshuffle_all benchmark_gemm_multi_d_all && \ + python3 ../tile_engine/ops/gemm/gemm_universal/gemm_benchmark.py . --problem-sizes "1024,1024,1024" --warmup 5 --repeat 5 --verbose --json results.json && \ + python3 ../tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.py . --problem-sizes "1024,1024,1024" --warmup 5 --repeat 5 --verbose --json results.json && \ + python3 ../tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.py . --problem-sizes "1024,1024,1024" --warmup 5 --repeat 5 --verbose --json results.json """ } steps{ buildHipClangJobAndReboot(setup_args:setup_args, build_type: 'Release', execute_cmd: execute_args) From 841351afa4f44f2038c5b5bbe07490a168d6823a Mon Sep 17 00:00:00 2001 From: ThruptiRajLakshmanaGowda Date: Tue, 13 Jan 2026 18:00:45 +0000 Subject: [PATCH 4/6] Removing gfx90a from Tile Engine Run --- Jenkinsfile | 31 ------------------------------- 1 file changed, 31 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 91c537d8b92..a554562c249 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -1670,37 +1670,6 @@ pipeline { } parallel { - stage("Run TILE_ENGINE_GEMM Tests on gfx90a") - { - when { - beforeAgent true - expression { params.RUN_TILE_ENGINE_GEMM_TESTS.toBoolean() } - } - agent{ label rocmnode("gfx90a") } - 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="gfx90a" \ - -D GEMM_UNIVERSAL_DATATYPE="fp8;fp16" \ - -D GEMM_UNIVERSAL_LAYOUT="rcr;rrr;crr;ccr" \ - -D GEMM_STREAMK_DATATYPE="fp8;fp16" \ - -D GEMM_STREAMK_LAYOUT="rcr" \ - -D GEMM_MULTI_D_DATATYPE="fp16" \ - -D GEMM_MULTI_D_LAYOUT="rcrr;rrrr;crrr;ccrr" \ - -D GEMM_PRESHUFFLE_DATATYPE="fp16;fp8;bf16;bf8" \ - -D GEMM_PRESHUFFLE_LAYOUT="rcr" .. && \ - ninja -j${nthreads()} benchmark_gemm_universal_all benchmark_gemm_preshuffle_all benchmark_gemm_multi_d_all benchmark_gemm_streamk_all && \ - python3 ../tile_engine/ops/gemm/gemm_universal/gemm_benchmark.py . --problem-sizes "1024,1024,1024" --warmup 5 --repeat 5 --verbose --json results.json && \ - python3 ../tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.py . --problem-sizes "1024,1024,1024" --warmup 5 --repeat 5 --verbose --json results.json && \ - python3 ../tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.py . --problem-sizes "1024,1024,1024" --warmup 5 --repeat 5 --verbose --json results.json """ - } - steps{ - buildHipClangJobAndReboot(setup_args:setup_args, build_type: 'Release', execute_cmd: execute_args) - cleanWs() - } - } stage("Run TILE_ENGINE_GEMM Tests on gfx942") { when { From b5b79a89a9c495cfb9f08572d8e2a67282f3e918 Mon Sep 17 00:00:00 2001 From: ThruptiRajLakshmanaGowda Date: Tue, 13 Jan 2026 21:50:26 +0000 Subject: [PATCH 5/6] [CK TILE ENGINE] increasing ci configs for BASIC case --- .../ops/gemm/gemm_multi_d/configs/default_ci_config.json | 7 +++++-- .../gemm/gemm_preshuffle/configs/default_ci_config.json | 1 + .../ops/gemm/gemm_universal/configs/default_ci_config.json | 7 +++++-- 3 files changed, 11 insertions(+), 4 deletions(-) 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 index 2df04d0ac1e..0698786c337 100644 --- 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 @@ -49,7 +49,9 @@ "trait_config": { "pipeline": { "values": [ - "compv4" + "compv3", + "compv4", + "mem" ] }, "scheduler": { @@ -60,7 +62,8 @@ }, "epilogue": { "values": [ - "cshuffle" + "cshuffle", + "default" ] }, "pad_m": { 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 index b32d42dabfa..868debad3f2 100644 --- a/tile_engine/ops/gemm/gemm_preshuffle/configs/default_ci_config.json +++ b/tile_engine/ops/gemm/gemm_preshuffle/configs/default_ci_config.json @@ -59,6 +59,7 @@ }, "epilogue": { "values": [ + "default", "cshuffle" ] }, 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 index 2dd8230edcb..38376a410b0 100644 --- a/tile_engine/ops/gemm/gemm_universal/configs/default_ci_config.json +++ b/tile_engine/ops/gemm/gemm_universal/configs/default_ci_config.json @@ -49,7 +49,9 @@ "trait_config": { "pipeline": { "values": [ - "compv4" + "compv3", + "compv4", + "mem" ] }, "scheduler": { @@ -60,7 +62,8 @@ }, "epilogue": { "values": [ - "cshuffle" + "cshuffle", + "default" ] }, "pad_m": { From 32d78f1a141330023f792a06ff2335bfa2a05d4b Mon Sep 17 00:00:00 2001 From: ThruptiRajLakshmanaGowda Date: Tue, 13 Jan 2026 23:47:47 +0000 Subject: [PATCH 6/6] Setting RUN_TILE_ENGINE_BASIC_TESTS to ON by default --- Jenkinsfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index a554562c249..406a5644c93 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -1201,8 +1201,8 @@ pipeline { 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)") + defaultValue: true, + description: "Run the tile_engine_basic tests (default: ON)") booleanParam( name: "RUN_TILE_ENGINE_GEMM_TESTS", defaultValue: false,