From 797aff8d496ee745d94fa000a28e4405b4f33e0b Mon Sep 17 00:00:00 2001
From: kyle-256 <Kyle.Zhao@amd.com>
Date: Fri, 23 Jan 2026 08:59:50 +0000
Subject: [PATCH] update test codes

---
 example/ck_tile/03_gemm/universal_gemm.cpp    |  2 +-
 .../ck_tile/17_grouped_gemm/grouped_gemm.cpp  | 21 ++++++++-
 .../ck_tile/17_grouped_gemm/grouped_gemm.hpp  | 43 +++++++++++++++++++
 3 files changed, 64 insertions(+), 2 deletions(-)

diff --git a/example/ck_tile/03_gemm/universal_gemm.cpp b/example/ck_tile/03_gemm/universal_gemm.cpp
index 5f6c149539..14b33dbb5c 100644
--- a/example/ck_tile/03_gemm/universal_gemm.cpp
+++ b/example/ck_tile/03_gemm/universal_gemm.cpp
@@ -46,7 +46,7 @@ int main(int argc, char* argv[])
     try
     {   run_gemm_example<GemmConfigComputeV3_1>(arg_parser);
         run_gemm_example<GemmConfigComputeV4>(arg_parser);
-        run_gemm_example<GemmConfigComputeAsync>(arg_parser);
+        // run_gemm_example<GemmConfigComputeAsync>(arg_parser);
         return 0;
     }
     catch(const std::runtime_error& e)
diff --git a/example/ck_tile/17_grouped_gemm/grouped_gemm.cpp b/example/ck_tile/17_grouped_gemm/grouped_gemm.cpp
index 781e0c0932..60acddee5f 100644
--- a/example/ck_tile/17_grouped_gemm/grouped_gemm.cpp
+++ b/example/ck_tile/17_grouped_gemm/grouped_gemm.cpp
@@ -326,6 +326,12 @@ int run_grouped_gemm_example_with_n_check(int argc, char* argv[])
             return run_gemm_example_prec_type<GemmConfigComputeV3_256x256_w16_k1<ck_tile::bf16_t>,
                                               ck_tile::bf16_t>(a_layout, b_layout, argc, argv);
         }
+        else if(config == "compute_v3_256x256x32_1x4x1_16x16x32")
+        {
+            return run_gemm_example_prec_type<
+                GemmConfigComputeV3_256x256x32_1x4x1_16x16x32<ck_tile::bf16_t>,
+                ck_tile::bf16_t>(a_layout, b_layout, argc, argv);
+        }
         else if(config == "compute_v3_128x128x64_2x2x1_32x32x16")
         {
             return run_gemm_example_prec_type<
@@ -344,14 +350,27 @@ int run_grouped_gemm_example_with_n_check(int argc, char* argv[])
                 GemmConfigComputeV3_128x256x64_1x4x1_32x32x16<ck_tile::bf16_t>,
                 ck_tile::bf16_t>(a_layout, b_layout, argc, argv);
         }
+        else if(config == "compute_v4")
+        {
+            return run_gemm_example_prec_type<GemmConfigComputeV4<ck_tile::bf16_t>, ck_tile::bf16_t>(
+                a_layout, b_layout, argc, argv);
+        }
+        else if(config == "compute_v4_128x128x64_2x2x1_32x32x16")
+        {
+            return run_gemm_example_prec_type<
+                GemmConfigComputeV4_128x128x64_2x2x1_32x32x16<ck_tile::bf16_t>,
+                ck_tile::bf16_t>(a_layout, b_layout, argc, argv);
+        }
         else
         {
             throw std::runtime_error(
                 "Unknown config: " + config +
                 ". Use: compute_v3, compute_v3_128x128, compute_v3_256x256_w16_k1, "
+                "compute_v3_256x256x32_1x4x1_16x16x32, "
                 "compute_v3_128x128x64_2x2x1_32x32x16, "
                 "compute_v3_128x128x128_2x2x1_32x32x16, "
-                "compute_v3_128x256x64_1x4x1_32x32x16");
+                "compute_v3_128x256x64_1x4x1_32x32x16, "
+                "compute_v4, compute_v4_128x128x64_2x2x1_32x32x16");
         }
     }
     else
diff --git a/example/ck_tile/17_grouped_gemm/grouped_gemm.hpp b/example/ck_tile/17_grouped_gemm/grouped_gemm.hpp
index 6ad6648fca..71f4f35272 100644
--- a/example/ck_tile/17_grouped_gemm/grouped_gemm.hpp
+++ b/example/ck_tile/17_grouped_gemm/grouped_gemm.hpp
@@ -258,6 +258,27 @@ struct GemmConfigComputeV3_256x128 : public GemmConfigBase
     static constexpr int kBlockPerCu = 1;
 };
 
+// 256x256x32, warp 1x4x1, warp tile 16x16x32
+template <typename PrecType>
+struct GemmConfigComputeV3_256x256x32_1x4x1_16x16x32 : public GemmConfigBase
+{
+    static constexpr ck_tile::index_t M_Tile = 256;
+    static constexpr ck_tile::index_t N_Tile = 256;
+    static constexpr ck_tile::index_t K_Tile = 64 / sizeof(PrecType); // 32 for fp16/bf16
+
+    static constexpr ck_tile::index_t M_Warp = 1;
+    static constexpr ck_tile::index_t N_Warp = 4;
+    static constexpr ck_tile::index_t K_Warp = 1;
+
+    static constexpr ck_tile::index_t M_Warp_Tile = 16;
+    static constexpr ck_tile::index_t N_Warp_Tile = 16;
+    static constexpr ck_tile::index_t K_Warp_Tile = 32;
+
+    static constexpr bool DoubleSmemBuffer          = false;
+    static constexpr ck_tile::GemmPipeline Pipeline = ck_tile::GemmPipeline::COMPUTE_V3;
+    static constexpr int kBlockPerCu                = 1;
+};
+
 // 128x128x64, warp 2x2x1, warp tile 32x32x16
 template <typename PrecType>
 struct GemmConfigComputeV3_128x128x64_2x2x1_32x32x16 : public GemmConfigBase
@@ -372,6 +393,28 @@ struct GemmConfigComputeV4 : public GemmConfigBase
     static constexpr int kBlockPerCu = 2;
 };
 
+// 128x128x64, warp 2x2x1, warp tile 32x32x16 (V4)
+template <typename PrecType>
+struct GemmConfigComputeV4_128x128x64_2x2x1_32x32x16 : public GemmConfigBase
+{
+    static constexpr ck_tile::index_t M_Tile = 128;
+    static constexpr ck_tile::index_t N_Tile = 128;
+    static constexpr ck_tile::index_t K_Tile = 128 / sizeof(PrecType); // 64 for fp16/bf16
+
+    static constexpr ck_tile::index_t M_Warp = 2;
+    static constexpr ck_tile::index_t N_Warp = 2;
+    static constexpr ck_tile::index_t K_Warp = 1;
+
+    static constexpr ck_tile::index_t M_Warp_Tile = 32;
+    static constexpr ck_tile::index_t N_Warp_Tile = 32;
+    static constexpr ck_tile::index_t K_Warp_Tile = 16;
+
+    static constexpr bool DoubleSmemBuffer          = true;
+    static constexpr ck_tile::GemmPipeline Pipeline = ck_tile::GemmPipeline::COMPUTE_V4;
+
+    static constexpr int kBlockPerCu = 1;
+};
+
 template <typename PrecType>
 struct GemmConfigComputeV4_V2 : public GemmConfigBase
 {