diff --git a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf16.cpp b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf16.cpp index 77b3be8684..e1a64c8656 100644 --- a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf16.cpp +++ b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf16.cpp @@ -12,9 +12,8 @@ using GemmConfig = GemmConfigQuantPrefill; QuantGroupSize, \ ck_tile::QuantType::BQuantGrouped>(arg_parser); -void bquant_quantgrouped_mx_bf16bf16_instance_factory( - std::unordered_map>& lut) -{ +static auto _ = []() { + auto& lut = get_kernel_lut(); using TypeConfig = decltype(GemmQuantTypeConfig>; return RUN_GEMM_EXAMPLE_PREC_TYPE; }; -} + return 0; +}(); diff --git a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf8.cpp b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf8.cpp index 9ac00b62dd..0eb2a0ce34 100644 --- a/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf8.cpp +++ b/example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf8.cpp @@ -11,9 +11,8 @@ using GemmConfig = GemmConfigMixedPrecision; QuantGroupSize, \ ck_tile::QuantType::BQuantGrouped>(arg_parser); -void bquant_quantgrouped_mx_bf16bf8_instance_factory( - std::unordered_map>& lut) -{ +static auto _ = []() { + auto& lut = get_kernel_lut(); using TypeConfig = decltype(GemmQuantTypeConfig>; return RUN_GEMM_EXAMPLE_PREC_TYPE; }; -} + return 0; +}(); diff --git a/example/ck_tile/38_block_scale_gemm/gemm_quant.cpp b/example/ck_tile/38_block_scale_gemm/gemm_quant.cpp index d21c94ff89..dc4d1ad814 100644 --- a/example/ck_tile/38_block_scale_gemm/gemm_quant.cpp +++ b/example/ck_tile/38_block_scale_gemm/gemm_quant.cpp @@ -123,4 +123,5 @@ int main(int argc, char* argv[]) "group_size not supported." << std::endl; return -1; + } } diff --git a/include/ck_tile/ops/gemm_quant/block/block_universal_gemm_as_bs_bquant_cr.hpp b/include/ck_tile/ops/gemm_quant/block/block_universal_gemm_as_bs_bquant_cr.hpp index 277a249614..5545112a4f 100644 --- a/include/ck_tile/ops/gemm_quant/block/block_universal_gemm_as_bs_bquant_cr.hpp +++ b/include/ck_tile/ops/gemm_quant/block/block_universal_gemm_as_bs_bquant_cr.hpp @@ -284,9 +284,9 @@ struct BQuantBlockUniversalGemmAsBsCr static_for<0, Traits::QScalesPerBlockRow, 1>{}([&](auto kQScale) { // B scale register offset constexpr index_t reg_offset = [&]() { - if constexpr(GemmTraits::QuantGroupSize::kN >= (NWarp * WarpGemm::kN)) - return (nIter * NWarp * WarpGemm::kN) / GemmTraits::QuantGroupSize::kN * - Traits::KQPerBlock + + if constexpr(GemmTraits::BQuantGroupSize::kN >= (NWarp * WarpGemm::kN)) + return (nIter * NWarp * WarpGemm::kN) / + GemmTraits::BQuantGroupSize::kN * Traits::KQPerBlock + kQScale; else { diff --git a/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp b/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp index de92d45763..589ec44ab6 100644 --- a/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp +++ b/include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp @@ -33,9 +33,9 @@ struct MxGemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3 using BLDSType = std::conditional_t; - using BQDataType = remove_cvref_t; - using CDataType = remove_cvref_t; - using BlockGemmShape = remove_cvref_t; + using BQDataType = remove_cvref_t; + using CDataType = remove_cvref_t; + using BlockGemmShape = remove_cvref_t; using BQuantGroupSize = remove_cvref_t; static_assert(BQuantGroupSize::kM == 1, "only N/K blocks for BQuant kernel!");