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

[CK TILE ENGINE] CI fix for Basic Tile Engine (#3554)

Browse Source 
* memory op changes

* memory op changes

* Fixing TILE_ENGINE_BASIC in Tile Engine

* Removing gfx90a from Tile Engine Run

* [CK TILE ENGINE] increasing ci configs for BASIC case

* Setting RUN_TILE_ENGINE_BASIC_TESTS to ON by default

---------

Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>

[ROCm/composable_kernel commit: 51027474af]
This commit is contained in:
Thrupti Raj Lakshmana Gowda
2026-01-13 18:20:30 -06:00
committed by GitHub
parent f444eab66c
commit 183c01c8f1
5 changed files with 31 additions and 76 deletions
Download Patch File Download Diff File Expand all files Collapse all files

52
tile_engine/ops/gemm/gemm_instance_builder.py
View File

@@ -643,40 +643,31 @@ struct SelectedKernel {{
using GemmPipeline = {pipeline_impl_map.get(pipeline)}<UniversalGemmProblem>;"""
# 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;"""
constexpr auto scheduler = {scheduler_type_map.get(scheduler)};"""
# UniversalGemmProblem
if self.kernel_name_prefix in ["gemm_universal"]:
instance_code += """
using UniversalGemmProblem = ck_tile::UniversalGemmPipelineProblem<
ADataType,
BDataType,
AccDataType,
TileShape,
ck_tile::TileGemmUniversalTraits<kPadM, kPadN, kPadK, DoubleSmemBuffer,
ALayout, BLayout, CLayout, TransposeC,
UseStructuredSparsity, UsePersistentKernel,
NumWaveGroups, Preshuffle>,
scheduler>;"""
using UniversalGemmProblem = ck_tile::UniversalGemmPipelineProblem<
ADataType,
BDataType,
AccDataType,
TileShape,
ck_tile::TileGemmUniversalTraits<kPadM, kPadN, kPadK, DoubleSmemBuffer,
ALayout, BLayout, CLayout, TransposeC,
UseStructuredSparsity, UsePersistentKernel,
NumWaveGroups, Preshuffle>,
scheduler>;"""
# GemmPipeline
if self.kernel_name_prefix in ["gemm_universal"]:
instance_code += f"""
using GemmPipeline = {pipeline_impl_map.get(pipeline)}<UniversalGemmProblem>;"""
using GemmPipeline = {pipeline_impl_map.get(pipeline)}<UniversalGemmProblem>;"""
# Epilogue
instance_code += self.populate_epilogue(epilogue)
@@ -748,23 +739,8 @@ struct SelectedKernel {{
ck_tile::make_kernel<kBlockPerCu>(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<ck_tile::memory_operation_enum,
ck_tile::memory_operation_enum::set>{});
} else {
ave_time = Run(ck_tile::integral_constant<ck_tile::memory_operation_enum,
ck_tile::memory_operation_enum::atomic_add>{});
}
return ave_time;
}
};
}}
}};
"""
return instance_code

7
tile_engine/ops/gemm/gemm_multi_d/configs/default_ci_config.json
View File

@@ -49,7 +49,9 @@
"trait_config": {
"pipeline": {
"values": [
"compv4"
"compv3",
"compv4",
"mem"
]
},
"scheduler": {
@@ -60,7 +62,8 @@
},
"epilogue": {
"values": [
"cshuffle"
"cshuffle",
"default"
]
},
"pad_m": {

1
tile_engine/ops/gemm/gemm_preshuffle/configs/default_ci_config.json
View File

@@ -59,6 +59,7 @@
},
"epilogue": {
"values": [
"default",
"cshuffle"
]
},

7
tile_engine/ops/gemm/gemm_universal/configs/default_ci_config.json
View File

@@ -49,7 +49,9 @@
"trait_config": {
"pipeline": {
"values": [
"compv4"
"compv3",
"compv4",
"mem"
]
},
"scheduler": {
@@ -60,7 +62,8 @@
},
"epilogue": {
"values": [
"cshuffle"
"cshuffle",
"default"
]
},
"pad_m": {