From 82fcd3c0bd7be7a95db01297f778405319546f32 Mon Sep 17 00:00:00 2001
From: khuagarw <khuagarw@amd.com>
Date: Fri, 19 Dec 2025 01:49:22 +0000
Subject: [PATCH] disabling 128 on non preshufflequant

---
 .../gemm_bquant_quantgrouped_fp8.cpp          | 12 ++---
 ...mm_bquant_quantgrouped_preshufflequant.cpp | 54 +++++++++----------
 2 files changed, 32 insertions(+), 34 deletions(-)

diff --git a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_fp8.cpp b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_fp8.cpp
index 8ad5529d19..fc44869a1d 100644
--- a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_fp8.cpp
+++ b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_fp8.cpp
@@ -49,10 +49,10 @@ void bquant_quantgrouped_fp8_instance_factory(
             using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 64, 128>>;
             return RUN_GEMM_EXAMPLE_PREC_TYPE;
         };
-    lut[hash_multiple_strings(
-        {"fp8", "bquant", "non-preshuffleb", "non-preshufflequant", "1x128x128"})] =
-        [](const ck_tile::ArgParser& arg_parser) {
-            using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 128, 128>>;
-            return RUN_GEMM_EXAMPLE_PREC_TYPE;
-        };
+    // lut[hash_multiple_strings(
+    //     {"fp8", "bquant", "non-preshuffleb", "non-preshufflequant", "1x128x128"})] =
+    //     [](const ck_tile::ArgParser& arg_parser) {
+    //         using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 128, 128>>;
+    //         return RUN_GEMM_EXAMPLE_PREC_TYPE;
+    //     };
 }
diff --git a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant.cpp b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant.cpp
index e1c8ed3273..fd39b31a53 100644
--- a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant.cpp
+++ b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshufflequant.cpp
@@ -4,7 +4,7 @@
 #include "run_gemm_quant_example.inc"
 
 template <typename T>
-using GemmConfig = GemmConfigPreshuffleQuantDecode<T>;
+using GemmConfig = GemmConfigPreshuffleBQuantPrefill<T>; // GemmConfigPreshuffleQuantDecode<T>;
 // //GemmConfigPreshuffleBQuantPrefill<T>;
 
 void bquant_quantgrouped_preshufflequant_instance_factory(
@@ -24,33 +24,31 @@ void bquant_quantgrouped_preshufflequant_instance_factory(
                                               ck_tile::QuantType::BQuantGrouped>(arg_parser);
         };
 
-    // lut[hash_multiple_strings({"fp8", "bquant", "non-preshuffleb", "preshufflequant",
-    // "1x8x128"})] =
-    //     [](const ck_tile::ArgParser& arg_parser) {
-    //         using TypeConfig     = decltype(GemmQuantTypeConfig<ck_tile::fp8_t,
-    //                                                             ck_tile::fp8_t,
-    //                                                             ck_tile::half_t,
-    //                                                             float>{});
-    //         using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 8, 128>>;
-    //         return run_gemm_example_prec_type<GemmConfig<ck_tile::fp8_t>,
-    //                                           TypeConfig,
-    //                                           QuantGroupSize,
-    //                                           ck_tile::QuantType::BQuantGrouped>(arg_parser);
-    //     };
-    // lut[hash_multiple_strings({"fp8",
-    //                            "bquant",
-    //                            "non-preshuffleb",
-    //                            "preshufflequant",
-    //                            "1x16x128"})] = [](const ck_tile::ArgParser& arg_parser) {
-    //     using TypeConfig =
-    //         decltype(GemmQuantTypeConfig<ck_tile::fp8_t, ck_tile::fp8_t, ck_tile::half_t,
-    //         float>{});
-    //     using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 16, 128>>;
-    //     return run_gemm_example_prec_type<GemmConfig<ck_tile::fp8_t>,
-    //                                       TypeConfig,
-    //                                       QuantGroupSize,
-    //                                       ck_tile::QuantType::BQuantGrouped>(arg_parser);
-    // };
+    lut[hash_multiple_strings({"fp8", "bquant", "non-preshuffleb", "preshufflequant", "1x8x128"})] =
+        [](const ck_tile::ArgParser& arg_parser) {
+            using TypeConfig     = decltype(GemmQuantTypeConfig<ck_tile::fp8_t,
+                                                                ck_tile::fp8_t,
+                                                                ck_tile::half_t,
+                                                                float>{});
+            using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 8, 128>>;
+            return run_gemm_example_prec_type<GemmConfig<ck_tile::fp8_t>,
+                                              TypeConfig,
+                                              QuantGroupSize,
+                                              ck_tile::QuantType::BQuantGrouped>(arg_parser);
+        };
+    lut[hash_multiple_strings({"fp8",
+                               "bquant",
+                               "non-preshuffleb",
+                               "preshufflequant",
+                               "1x16x128"})] = [](const ck_tile::ArgParser& arg_parser) {
+        using TypeConfig =
+            decltype(GemmQuantTypeConfig<ck_tile::fp8_t, ck_tile::fp8_t, ck_tile::half_t, float>{});
+        using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 16, 128>>;
+        return run_gemm_example_prec_type<GemmConfig<ck_tile::fp8_t>,
+                                          TypeConfig,
+                                          QuantGroupSize,
+                                          ck_tile::QuantType::BQuantGrouped>(arg_parser);
+    };
     lut[hash_multiple_strings({"fp8",
                                "bquant",
                                "non-preshuffleb",