diff --git a/example/ck_tile/38_block_scale_gemm/run_gemm_quant_example.inc b/example/ck_tile/38_block_scale_gemm/run_gemm_quant_example.inc index 2ffe146a75..bce716079e 100644 --- a/example/ck_tile/38_block_scale_gemm/run_gemm_quant_example.inc +++ b/example/ck_tile/38_block_scale_gemm/run_gemm_quant_example.inc @@ -505,7 +505,8 @@ int run_gemm_example_with_layouts(const ck_tile::ArgParser& arg_parser, { ck_tile::FillUniformDistribution{-2.0f, 3.0f, fill_seed(gen)}(a_m_k); } - ck_tile::FillUniformDistribution{0, 0.5f, fill_seed(gen)}(*aq_tensor_ptr); + ck_tile::FillUniformDistribution{-2.0f, 2.0f, fill_seed(gen)}( + *aq_tensor_ptr); ck_tile::FillUniformDistribution{-5.0f, 5.0f, fill_seed(gen)}(b_k_n); } else @@ -528,7 +529,7 @@ int run_gemm_example_with_layouts(const ck_tile::ArgParser& arg_parser, if constexpr(QuantMode == ck_tile::QuantType::BQuantGrouped) { ck_tile::FillConstant{static_cast(0x38)}(a_m_k); - ck_tile::FillConstant{static_cast(0x38)}(b_k_n); + ck_tile::FillConstant{static_cast(0x22)}(b_k_n); ck_tile::FillConstant{static_cast(0.5f)}(*bq_tensor_ptr); } else diff --git a/include/ck_tile/ops/gemm_quant/block/block_universal_gemm_ar_flatbr_bquant_cr.hpp b/include/ck_tile/ops/gemm_quant/block/block_universal_gemm_ar_flatbr_bquant_cr.hpp index 2e4f60d0b2..a80583f999 100644 --- a/include/ck_tile/ops/gemm_quant/block/block_universal_gemm_ar_flatbr_bquant_cr.hpp +++ b/include/ck_tile/ops/gemm_quant/block/block_universal_gemm_ar_flatbr_bquant_cr.hpp @@ -190,6 +190,7 @@ struct BlockGemmWeightPreshuffleBQuantARegBRegCReg scale_reg_dword = static_cast(scale_reg); } + // cross lane ops to get the value of scale_reg. int gathered_scale_reg = __builtin_amdgcn_ds_bpermute( pull_from_lane << 2, __builtin_bit_cast(int, scale_reg_dword)); diff --git a/include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp b/include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp index 8eb1ee4496..32f1279e93 100644 --- a/include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp +++ b/include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp @@ -390,7 +390,6 @@ struct QuantGroupedGemmKernel const auto& c_block_tile = GemmPipeline{}.template operator()(a_block_window, b_block_window, bq_block_window, - kargs.N, num_loop, tail_num, smem_ptr_0, @@ -458,7 +457,6 @@ struct QuantGroupedGemmKernel const auto& c_block_tile = GemmPipeline{}.template operator()(a_block_window, b_block_window, bq_block_window, - kargs.N, num_loop, has_hot_loop, tail_num, diff --git a/include/ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp b/include/ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp index 611ba21760..b92c9ee1fd 100644 --- a/include/ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp +++ b/include/ck_tile/ops/gemm_quant/pipeline/gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp @@ -70,7 +70,7 @@ struct WPQuantBPipelineAgBgCrV2 : public WeightPreshufflePipelineAGmemBGmemCRegV static constexpr bool PreshuffleQuant = Problem::Traits::PreshuffleQuant; static constexpr index_t KPerBlockBQ = - integer_divide_ceil(BlockGemmShape::kK, QuantGroupSize::kK); // 256/128 + integer_divide_ceil(BlockGemmShape::kK, QuantGroupSize::kK); static constexpr index_t QScalesPerBlockRow = integer_divide_ceil(kKPerBlock, QuantGroupSize::kK);