ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 01:27:34 +00:00
Code Issues Packages Projects Releases Wiki Activity

Renaming A/B (lds) load distribution as they have different semantic + cleanup

Browse Source
This commit is contained in:
Damien Lejeune
2025-07-28 15:04:39 +00:00
parent 6044797613
commit 00b153b99c
3 changed files with 8 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
include/ck_tile/ops/gemm/block/block_universal_gemm_ar_br_cr.hpp
View File

@@ -39,9 +39,6 @@ struct BlockUniversalGemmArBrCr : public BlockUniversalGemmBase<Problem_, Policy
using GemmTraits = typename Base::template GemmTraits_<Problem_, Policy_>;
public:
using Base::MakeABlockDistributionEncode;
using Base::MakeBBlockDistributionEncode;
using Base::MakeCBlockTile;
using typename Base::WarpGemm;
private:

1
include/ck_tile/ops/gemm/block/block_universal_gemm_ar_bs_cr.hpp
View File

@@ -40,7 +40,6 @@ struct BlockUniversalGemmArBsCr : public BlockUniversalGemmBase<Problem_, Policy
using GemmTraits = typename Base::template GemmTraits_<Problem_, Policy_>;
public:
using Base::MakeABlockDistributionEncode;
using Base::MakeBBlockDistributionEncode;
using Base::MakeCBlockTile;
using typename Base::WarpGemm;

16
include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_base.hpp
View File

@@ -71,10 +71,10 @@ struct GemmPipelineAgBgCrImplBase
return make_tuple(std::move(a_lds_block), std::move(b_lds_block));
}
template <typename ADramBlockWindow, typename ALdsTensorView, typename ALdsLoadTileDistr>
template <typename ADramBlockWindow, typename ALdsTensorView, typename ALoadTileDistr>
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<ALayout, tensor_layout::gemm::ColumnMajor>;
@@ -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<MPerBlock>{}, number<KPerBlock>{}),
{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 <typename BDramBlockWindow, typename BLdsTensorView, typename BLdsLoadTileDistr>
template <typename BDramBlockWindow, typename BLdsTensorView, typename BLoadTileDistr>
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<BLayout, tensor_layout::gemm::RowMajor>;
@@ -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<NPerBlock>{}, number<KPerBlock>{}),
{0, 0},
BLdsLoadTileDistr{});
BLoadTileDistr{});
return make_tuple(std::move(b_copy_dram_window),
std::move(b_copy_lds_window),