ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 09:08:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix rebase issues

Browse Source
This commit is contained in:
Enrico Degregori
2026-01-30 13:12:37 +00:00
parent b6a66c19e8
commit 3dde9eb13b
5 changed files with 15 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf16.cpp
View File

@@ -12,9 +12,8 @@ using GemmConfig = GemmConfigQuantPrefill<T>;
QuantGroupSize, \
ck_tile::QuantType::BQuantGrouped>(arg_parser);
void bquant_quantgrouped_mx_bf16bf16_instance_factory(
std::unordered_map<size_t, std::function<int(const ck_tile::ArgParser&)>>& lut)
{
static auto _ = []() {
auto& lut = get_kernel_lut();
using TypeConfig = decltype(GemmQuantTypeConfig<ck_tile::bf16_t,
ck_tile::bf16_t,
ck_tile::bf16_t,
@@ -32,4 +31,5 @@ void bquant_quantgrouped_mx_bf16bf16_instance_factory(
using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 1, 64>>;
return RUN_GEMM_EXAMPLE_PREC_TYPE;
};
}
return 0;
}();

8
example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_mx_bf16bf8.cpp
View File

@@ -11,9 +11,8 @@ using GemmConfig = GemmConfigMixedPrecision;
QuantGroupSize, \
ck_tile::QuantType::BQuantGrouped>(arg_parser);
void bquant_quantgrouped_mx_bf16bf8_instance_factory(
std::unordered_map<size_t, std::function<int(const ck_tile::ArgParser&)>>& lut)
{
static auto _ = []() {
auto& lut = get_kernel_lut();
using TypeConfig = decltype(GemmQuantTypeConfig<ck_tile::bf16_t,
ck_tile::bf8_t,
ck_tile::bf16_t,
@@ -31,4 +30,5 @@ void bquant_quantgrouped_mx_bf16bf8_instance_factory(
using QuantGroupSize = ck_tile::QuantGroupShape<ck_tile::sequence<1, 1, 64>>;
return RUN_GEMM_EXAMPLE_PREC_TYPE;
};
}
return 0;
}();

1
example/ck_tile/38_block_scale_gemm/gemm_quant.cpp
View File

@@ -123,4 +123,5 @@ int main(int argc, char* argv[])
"group_size not supported."
<< std::endl;
return -1;
}
}

6
include/ck_tile/ops/gemm_quant/block/block_universal_gemm_as_bs_bquant_cr.hpp
View File

@@ -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
{

6
include/ck_tile/ops/gemm_quant/pipeline/gemm_mx_pipeline_ag_bg_cr_v3.hpp
View File

@@ -33,9 +33,9 @@ struct MxGemmPipelineAgBgCrCompV3 : public BaseGemmPipelineAgBgCrCompV3<Problem>
using BLDSType = std::conditional_t<IsCastBeforeLDS, BDqDataType, BDataType>;
using BQDataType = remove_cvref_t<typename Problem::BQDataType>;
using CDataType = remove_cvref_t<typename Problem::CDataType>;
using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>;
using BQDataType = remove_cvref_t<typename Problem::BQDataType>;
using CDataType = remove_cvref_t<typename Problem::CDataType>;
using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>;
using BQuantGroupSize = remove_cvref_t<typename Problem::BQuantGroupSize>;
static_assert(BQuantGroupSize::kM == 1, "only N/K blocks for BQuant kernel!");