From 2d86cd0081ca8eea80d1ea364e3e710b8675e323 Mon Sep 17 00:00:00 2001 From: Sami Remes Date: Mon, 27 Oct 2025 14:28:06 +0000 Subject: [PATCH] fix formatting --- .../block_universal_gemm_as_bs_bquant_cr.hpp | 6 ++-- .../pipeline/gemm_group_quant_utils.hpp | 29 ++++++++++--------- 2 files changed, 20 insertions(+), 15 deletions(-) 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 8017a82bc6..56b34eff93 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 @@ -345,12 +345,14 @@ struct BQuantBlockUniversalGemmAsBsCr : public BlockGemmBQuantBase constexpr index_t reg_offset = [&]() { if constexpr(Traits::NQPerBlock >= Traits::NIterPerWarp) { - // Each nIter and warp/thread has its own scale - tile dstr handles the proper loading + // Each nIter and warp/thread has its own scale - tile dstr handles + // the proper loading return nIter * Traits::BQPerBlock + kQScale; } else { - // Many N warps/iters share the same scale, index from full [NQPerBlock=1, BQPerBlock] matrix + // Many N warps/iters share the same scale, index from full + // [NQPerBlock=1, BQPerBlock] matrix static_assert(Traits::NQPerBlock == 1); return kQScale; } diff --git a/include/ck_tile/ops/gemm_quant/pipeline/gemm_group_quant_utils.hpp b/include/ck_tile/ops/gemm_quant/pipeline/gemm_group_quant_utils.hpp index 2dac5cddb1..23256242b3 100644 --- a/include/ck_tile/ops/gemm_quant/pipeline/gemm_group_quant_utils.hpp +++ b/include/ck_tile/ops/gemm_quant/pipeline/gemm_group_quant_utils.hpp @@ -216,12 +216,13 @@ struct tile_distribution_encoding_pattern_bq : public tile_distribution_encoding constexpr index_t XR = get_warp_size() / NQPerIter; static_assert(YPerTile == NQPerIter * NWarps * NIterPerWarp); return make_static_tile_distribution( - tile_distribution_encoding, - tuple, sequence>, - tuple, sequence<0, 1>>, - tuple, sequence<1, 2>>, - sequence<1, 2>, - sequence<0, 0>>{}); + tile_distribution_encoding< + sequence, + tuple, sequence>, + tuple, sequence<0, 1>>, + tuple, sequence<1, 2>>, + sequence<1, 2>, + sequence<0, 0>>{}); } else if constexpr(YPerTile >= NIterPerWarp) { @@ -230,16 +231,18 @@ struct tile_distribution_encoding_pattern_bq : public tile_distribution_encoding constexpr index_t XR = get_warp_size() / NQPerIter; static_assert(YPerTile == NQPerIter * NIterPerWarp); return make_static_tile_distribution( - tile_distribution_encoding, - tuple, sequence>, - tuple, sequence<0, 1>>, - tuple, sequence<2, 1>>, - sequence<1, 2>, - sequence<0, 0>>{}); + tile_distribution_encoding< + sequence, + tuple, sequence>, + tuple, sequence<0, 1>>, + tuple, sequence<2, 1>>, + sequence<1, 2>, + sequence<0, 0>>{}); } else { - // larger NQ block size, multiple iters/warps use same scales -> replicate to all threads + // larger NQ block size, multiple iters/warps use same scales -> replicate to all + // threads return make_static_tile_distribution( tile_distribution_encoding, tuple, sequence>,