From 55a1a331dc2c1da425768ad832ad0bfc68032464 Mon Sep 17 00:00:00 2001
From: Philip Maybank <pmaybank@amd.com>
Date: Tue, 12 Aug 2025 17:46:21 +0100
Subject: [PATCH] debugging warp-level gemm classes

---
 include/ck_tile/core/arch/arch.hpp                    | 11 +++++++++++
 ...gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp |  2 +-
 2 files changed, 12 insertions(+), 1 deletion(-)

diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp
index 34ce76244a..86de1b8787 100644
--- a/include/ck_tile/core/arch/arch.hpp
+++ b/include/ck_tile/core/arch/arch.hpp
@@ -320,6 +320,17 @@ struct GfxConfig
             return Gfx90x::kCM0PerLane;
         }        
     }    
+    static constexpr index_t get_k_cm1_per_lane()
+    {
+        if constexpr (GfxId == 1200)
+        {
+            return Gfx120x::kCM1PerLane;
+        }
+        else if constexpr (GfxId == 900)
+        {
+            return Gfx90x::kCM1PerLane;
+        }        
+    }        
 };
 
 } // namespace ck_tile
diff --git a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp
index 290c64c741..6a52b9d3eb 100644
--- a/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp
+++ b/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_generic_impl_F16F16F32M16N16K16.hpp
@@ -91,7 +91,7 @@ struct WarpGemmAttributeGenericImplF16F16F32M16N16K16
     static constexpr index_t kCMLane     = GfxConfig<GfxId>::get_k_cm_lane();
     static constexpr index_t kCNLane     = 16;
     static constexpr index_t kCM0PerLane = GfxConfig<GfxId>::get_k_cm0_per_lane();
-    static constexpr index_t kCM1PerLane = 4;
+    static constexpr index_t kCM1PerLane = GfxConfig<GfxId>::get_k_cm1_per_lane();
 
     // c_vec += a_vec * b_vec
     template <bool post_nop_ = false>