From 4740aebc2b1771ebe6d644e1eafaa4e574ab1f7d Mon Sep 17 00:00:00 2001 From: Damien Lejeune Date: Thu, 31 Jul 2025 08:28:09 +0000 Subject: [PATCH] Undo renaming A/B (lds) load distribution for skip lds --- .../pipeline/gemm_pipeline_ag_bg_cr_base.hpp | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_base.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_base.hpp index ec40f8f32c..ad1624a99b 100644 --- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_base.hpp +++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_base.hpp @@ -91,10 +91,10 @@ struct GemmPipelineAgBgCrImplBase return make_tuple(std::move(a_lds_block), std::move(b_lds_block)); } - template + template CK_TILE_DEVICE constexpr auto GetAWindows(const ADramBlockWindow& a_dram_block_window, const ALdsTensorView& a_lds_block_view, - const ALoadTileDistr&, + const ALdsLoadTileDistr&, const array& offset = {0, 0}) const { constexpr bool is_col_major = std::is_same_v; @@ -116,7 +116,7 @@ struct GemmPipelineAgBgCrImplBase return make_tile_window(a_dram_block_window.get_bottom_tensor_view(), make_tuple(YPerTile{}, XPerTile{}), a_dram_block_window.get_window_origin() + offset, - ALoadTileDistr{}); + ALdsLoadTileDistr{}); } }(); @@ -133,10 +133,10 @@ struct GemmPipelineAgBgCrImplBase if constexpr(is_a_load_tr) return make_static_tile_distribution( typename InputTileDistributionTraits< - typename ALoadTileDistr::DstrEncode, + typename ALdsLoadTileDistr::DstrEncode, typename Problem::ADataType>::TransposedDstrEncode{}); else - return ALoadTileDistr{}; + return ALdsLoadTileDistr{}; }(); auto a_lds_gemm_window = make_tile_window(a_lds_block_view, a_lds_shape, {0, 0}, a_lds_load_tile_distr); @@ -146,10 +146,10 @@ struct GemmPipelineAgBgCrImplBase std::move(a_lds_gemm_window)); } - template + template CK_TILE_DEVICE constexpr auto GetBWindows(const BDramBlockWindow& b_dram_block_window, const BLdsTensorView& b_lds_block_view, - const BLoadTileDistr&, + const BLdsLoadTileDistr&, const array& offset = {0, 0}) const { constexpr bool is_row_major = std::is_same_v; @@ -170,7 +170,7 @@ struct GemmPipelineAgBgCrImplBase return make_tile_window(b_dram_block_window.get_bottom_tensor_view(), make_tuple(YPerTile{}, XPerTile{}), b_dram_block_window.get_window_origin() + offset, - BLoadTileDistr{}); + BLdsLoadTileDistr{}); } }(); @@ -189,10 +189,10 @@ struct GemmPipelineAgBgCrImplBase if constexpr(is_b_load_tr) return make_static_tile_distribution( typename InputTileDistributionTraits< - typename BLoadTileDistr::DstrEncode, + typename BLdsLoadTileDistr::DstrEncode, typename Problem::BDataType>::TransposedDstrEncode{}); else - return BLoadTileDistr{}; + return BLdsLoadTileDistr{}; }(); auto b_lds_gemm_window = make_tile_window(b_lds_block_view, b_lds_shape, {0, 0}, b_lds_load_tile_distr);