ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-08 08:07:06 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add static check verifying a/b transpose feature support

Browse Source
This commit is contained in:
Damien Lejeune
2025-09-04 08:29:39 +00:00
parent b0dd0b5234
commit 1fee5fd6e5
2 changed files with 24 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

29
include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp
View File

@@ -677,26 +677,43 @@ struct UniversalGemmPipelineAgBgCrPolicy
Problem::UseStructuredSparsity,
wg_attr_num_access>;
using BlockGemmPolicy = BlockGemmASmemBSmemCRegV1CustomPolicy<typename Problem::ADataType,
typename Problem::BDataType,
typename Problem::CDataType,
BlockWarps,
WarpGemm>;
if constexpr(Problem::SkipALds == false && Problem::SkipBLds == false)
{
using BlockGemmPolicy =
BlockGemmASmemBSmemCRegV1CustomPolicy<typename Problem::ADataType,
typename Problem::BDataType,
typename Problem::CDataType,
BlockWarps,
WarpGemm>;
return BlockUniversalGemmAsBsCr<Problem, BlockGemmPolicy>{};
}
else if constexpr(Problem::SkipALds == false && Problem::SkipBLds == true)
{
using BlockGemmPolicy =
BlockGemmASmemBRegCRegV1CustomPolicy<typename Problem::ADataType,
typename Problem::BDataType,
typename Problem::CDataType,
BlockWarps,
WarpGemm>;
return BlockUniversalGemmAsBrCr<Problem, BlockGemmPolicy>{};
}
else if constexpr(Problem::SkipALds == true && Problem::SkipBLds == false)
{
using BlockGemmPolicy =
BlockGemmARegBSmemCRegV1CustomPolicy<typename Problem::ADataType,
typename Problem::BDataType,
typename Problem::CDataType,
BlockWarps,
WarpGemm>;
return BlockUniversalGemmArBsCr<Problem, BlockGemmPolicy>{};
}
else
{
using BlockGemmPolicy = BlockGemmARegBRegCRegV1CustomPolicy<typename Problem::ADataType,
typename Problem::BDataType,
typename Problem::CDataType,
BlockWarps,
WarpGemm>;
return BlockUniversalGemmArBrCr<Problem, BlockGemmPolicy>{};
}
}