From 1fee5fd6e56eb0367687ab40ae262eebb12d4e08 Mon Sep 17 00:00:00 2001
From: Damien Lejeune <damien.lejeune@amd.com>
Date: Thu, 4 Sep 2025 08:29:39 +0000
Subject: [PATCH] Add static check verifying a/b transpose feature support

---
 ...emm_universal_pipeline_ag_bg_cr_policy.hpp | 29 +++++++++++++++----
 .../gemm/test_gemm_pipeline_kernel_types.hpp  |  2 +-
 2 files changed, 24 insertions(+), 7 deletions(-)

diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp
index 5937393f8c..1af40fbd4d 100644
--- a/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp
+++ b/include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp
@@ -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>{};
         }
     }
diff --git a/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp b/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp
index 7c561e994a..ad2f288882 100644
--- a/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp
+++ b/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp
@@ -57,7 +57,7 @@ using KernelTypesMem = ::testing::Types<
     std::tuple<    Col,     Col,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I64,        I32,        I32,        I16, Interwave,         Mem,  std::false_type, std::false_type>,
     std::tuple<    Col,     Col,     Row,       F8,        F8,          F32,       F16,        I256,        I256,         I64,        I32,        I32,        I16, Intrawave,         Mem,  std::false_type, std::false_type>,
     std::tuple<    Col,     Col,     Row,       F8,        F8,          F32,       F16,        I256,        I256,         I64,        I32,        I32,        I16, Interwave,         Mem,  std::false_type, std::false_type>,
-        // SkipALds/SkipBLds
+    // SkipALds/SkipBLds
     std::tuple<    Row,     Col,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I64,        I32,        I32,        I16, Interwave,         Mem,  std::false_type, std::true_type>,
     std::tuple<    Row,     Col,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I64,        I32,        I32,        I16, Interwave,         Mem,  std::true_type,  std::false_type>,
     std::tuple<    Row,     Col,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I64,        I32,        I32,        I16, Interwave,         Mem,  std::true_type,  std::true_type>,