From 8965f337ca1562efd58b4e817ad216e2e30402a3 Mon Sep 17 00:00:00 2001 From: "assistant-librarian[bot]" Date: Tue, 25 Nov 2025 21:12:15 +0000 Subject: [PATCH] Merge commit '37ea1600888f515e5dfb7153b75b2f06474d880d' into develop --- .../gemm_bquant_quantgrouped_preshuffleb.cpp | 5 +++++ ...uantgrouped_preshuffleb_preshufflequant.cpp | 5 +++++ .../ck_tile/38_block_scale_gemm/gemm_utils.hpp | 18 ++++++++++++++++++ 3 files changed, 28 insertions(+) diff --git a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb.cpp b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb.cpp index 898316fa6b..8271559cdc 100644 --- a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb.cpp +++ b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb.cpp @@ -3,8 +3,13 @@ #include "run_gemm_quant_example.inc" +#if CK_TILE_USE_WMMA +template +using GemmConfig = GemmConfigPreshuffleB_BQuant_Prefill_Wmma; +#else template using GemmConfig = GemmConfigPreshuffleB_BQuant_Prefill; +#endif void bquant_quantgrouped_preshuffleb_instance_factory( std::unordered_map>& lut) diff --git a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant.cpp b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant.cpp index e26edb6501..f89591c791 100644 --- a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant.cpp +++ b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant.cpp @@ -3,8 +3,13 @@ #include "run_gemm_quant_example.inc" +#if CK_TILE_USE_WMMA +template +using GemmConfig = GemmConfigPreshuffleB_PreshuffleBQuant_Prefill_Wmma; +#else template using GemmConfig = GemmConfigPreshuffleB_PreshuffleBQuant_Prefill; +#endif void bquant_quantgrouped_preshuffleb_preshufflequant_instance_factory( std::unordered_map>& lut) diff --git a/example/ck_tile/38_block_scale_gemm/gemm_utils.hpp b/example/ck_tile/38_block_scale_gemm/gemm_utils.hpp index 00118e4725..95b0a73ede 100644 --- a/example/ck_tile/38_block_scale_gemm/gemm_utils.hpp +++ b/example/ck_tile/38_block_scale_gemm/gemm_utils.hpp @@ -250,6 +250,24 @@ struct GemmConfigBQuantPrefill_Wmma : public GemmConfigBQuantPrefill static constexpr ck_tile::index_t K_Warp_Tile = 16; }; +template +struct GemmConfigPreshuffleB_BQuant_Prefill_Wmma + : public GemmConfigPreshuffleB_BQuant_Prefill +{ + 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 = 16; +}; + +template +struct GemmConfigPreshuffleB_PreshuffleBQuant_Prefill_Wmma + : public GemmConfigPreshuffleB_PreshuffleBQuant_Prefill +{ + 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 = 16; +}; + template