diff --git a/include/ck_tile/ops/gemm/block/block_universal_gemm_ar_br_cr.hpp b/include/ck_tile/ops/gemm/block/block_universal_gemm_ar_br_cr.hpp index 55593e44ad..0ca752b5be 100644 --- a/include/ck_tile/ops/gemm/block/block_universal_gemm_ar_br_cr.hpp +++ b/include/ck_tile/ops/gemm/block/block_universal_gemm_ar_br_cr.hpp @@ -39,9 +39,6 @@ struct BlockUniversalGemmArBrCr : public BlockUniversalGemmBase; public: - using Base::MakeABlockDistributionEncode; - using Base::MakeBBlockDistributionEncode; - using Base::MakeCBlockTile; using typename Base::WarpGemm; private: diff --git a/include/ck_tile/ops/gemm/block/block_universal_gemm_ar_bs_cr.hpp b/include/ck_tile/ops/gemm/block/block_universal_gemm_ar_bs_cr.hpp index fce137dffd..020de35471 100644 --- a/include/ck_tile/ops/gemm/block/block_universal_gemm_ar_bs_cr.hpp +++ b/include/ck_tile/ops/gemm/block/block_universal_gemm_ar_bs_cr.hpp @@ -40,7 +40,6 @@ struct BlockUniversalGemmArBsCr : public BlockUniversalGemmBase; public: - using Base::MakeABlockDistributionEncode; using Base::MakeBBlockDistributionEncode; using Base::MakeCBlockTile; using typename Base::WarpGemm; 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 1bed55c456..8eb5b167a1 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 @@ -71,10 +71,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 ALdsLoadTileDistr&) const + const ALoadTileDistr&) const { constexpr bool is_col_major = std::is_same_v; @@ -95,7 +95,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(), - ALdsLoadTileDistr{}); + ALoadTileDistr{}); } }(); @@ -107,17 +107,17 @@ struct GemmPipelineAgBgCrImplBase make_tile_window(a_lds_block_view, make_tuple(number{}, number{}), {0, 0}, - ALdsLoadTileDistr{}); + ALoadTileDistr{}); return make_tuple(std::move(a_copy_dram_window), std::move(a_copy_lds_window), 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 BLdsLoadTileDistr&) const + const BLoadTileDistr&) const { constexpr bool is_row_major = std::is_same_v; @@ -137,7 +137,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(), - BLdsLoadTileDistr{}); + BLoadTileDistr{}); } }(); @@ -151,7 +151,7 @@ struct GemmPipelineAgBgCrImplBase make_tile_window(b_lds_block_view, make_tuple(number{}, number{}), {0, 0}, - BLdsLoadTileDistr{}); + BLoadTileDistr{}); return make_tuple(std::move(b_copy_dram_window), std::move(b_copy_lds_window),