ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-25 07:14:37 +00:00
Code Issues Packages Projects Releases Wiki Activity

Enable padding blockscale for abquant (#3453)

Browse Source 
* Enable padding blockscale for abquant

* run clang-format

* Reduce unnecessary testing

* remove cout

[ROCm/composable_kernel commit: 7f68f3c4fa]
This commit is contained in:
kensclin
2025-12-25 01:12:40 +08:00
committed by GitHub
parent a2402950de
commit 0eb5d4a93f
8 changed files with 128 additions and 138 deletions
Download Patch File Download Diff File Expand all files Collapse all files

26
include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp
View File

@@ -412,32 +412,6 @@ struct QuantGemmKernel
return false;
}
if constexpr(kQuantType == QuantType::AQuantGrouped ||
kQuantType == QuantType::ABQuantGrouped)
{
if(kargs.QK_A % GemmPipeline::GetVectorSizeAQ() != 0)
{
if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
CK_TILE_ERROR("K_A is not a multiple of vector load size for A tensor!");
}
return false;
}
}
if constexpr(kQuantType == QuantType::BQuantGrouped ||
kQuantType == QuantType::ABQuantGrouped)
{
if(kargs.QK_B % GemmPipeline::GetVectorSizeBQ() != 0)
{
if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
CK_TILE_ERROR("K_B is not a multiple of vector load size for B tensor!");
}
return false;
}
}
if constexpr(std::is_same_v<ALayout, tensor_layout::gemm::RowMajor>)
{
if(kargs.K % (TilePartitioner::KPerBlock * kargs.k_batch) != 0 &&