From e0e29554ab7b8a26474725c4bf0c6c003ca34160 Mon Sep 17 00:00:00 2001 From: Jakub Piasecki Date: Mon, 5 May 2025 18:05:06 +0000 Subject: [PATCH] more fixes --- example/ck_tile/03_gemm/gemm_utils.hpp | 6 +- example/ck_tile/03_gemm/universal_gemm.cpp | 4 +- .../pipeline/gemm_pipeline_ag_bg_cr_base.hpp | 84 ++++--------------- test/ck_tile/gemm/test_gemm_pipeline_util.hpp | 2 - 4 files changed, 18 insertions(+), 78 deletions(-) diff --git a/example/ck_tile/03_gemm/gemm_utils.hpp b/example/ck_tile/03_gemm/gemm_utils.hpp index 88894b5d8f..0b4361d3b2 100644 --- a/example/ck_tile/03_gemm/gemm_utils.hpp +++ b/example/ck_tile/03_gemm/gemm_utils.hpp @@ -16,7 +16,7 @@ #define CK_TILE_PIPELINE_COMPUTE_V4 3 #ifndef CK_TILE_PIPELINE_DEFAULT -#define CK_TILE_PIPELINE_DEFAULT CK_TILE_PIPELINE_MEMORY +#define CK_TILE_PIPELINE_DEFAULT CK_TILE_PIPELINE_COMPUTE_V3 #endif #if(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_MEMORY) @@ -53,8 +53,8 @@ struct GemmConfig static constexpr bool DoubleSmemBuffer = false; - static constexpr bool SkipALds = false; - static constexpr bool SkipBLds = false; + static constexpr bool SkipALds = true; + static constexpr bool SkipBLds = true; #endif #if(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_COMPUTE_V3) // Compute friendly for Intrawave scheduler diff --git a/example/ck_tile/03_gemm/universal_gemm.cpp b/example/ck_tile/03_gemm/universal_gemm.cpp index 632404a9c8..3393a29a5d 100644 --- a/example/ck_tile/03_gemm/universal_gemm.cpp +++ b/example/ck_tile/03_gemm/universal_gemm.cpp @@ -144,9 +144,7 @@ float gemm_calc(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& << "\n File: " << __FILE__ << ":" << __LINE__ << ", in function: " << __func__; throw std::runtime_error(err.str()); } -#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_MEMORY || \ - CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_MEMORY_SKIP_B_LDS || \ - CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_MEMORY_SKIP_A_LDS) +#elif(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_MEMORY) // Tail pipeline One to Seven if(tail_num == ck_tile::TailNumber::One) { 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 816aefac12..1bed55c456 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 @@ -37,8 +37,8 @@ struct GemmPipelineAgBgCrImplBase const SrcBlockTile& src_block_tile, const ElementFunction& element_func) const { - const auto block_tile_tmp = tile_elementwise_in(element_func, src_block_tile); - store_tile(lds_tile_window, block_tile_tmp); + const auto block_tile = tile_elementwise_in(element_func, src_block_tile); + store_tile(lds_tile_window, block_tile); } template @@ -48,26 +48,6 @@ struct GemmPipelineAgBgCrImplBase load_tile(dst_block_tile, lds_tile_window); } - CK_TILE_DEVICE auto GetALdsTensorView(void* p_smem) const - { - // A tile in LDS - ADataType* __restrict__ p_a_lds = static_cast(p_smem); - constexpr auto a_lds_block_desc = Policy::template MakeALdsBlockDescriptor(); - auto a_lds_block = make_tensor_view(p_a_lds, a_lds_block_desc); - - return a_lds_block; - } - - CK_TILE_DEVICE auto GetBLdsTensorView(void* p_smem) const - { - // B tile in LDS - BDataType* __restrict__ p_b_lds = static_cast(p_smem); - constexpr auto b_lds_block_desc = Policy::template MakeBLdsBlockDescriptor(); - auto b_lds_block = make_tensor_view(p_b_lds, b_lds_block_desc); - - return b_lds_block; - } - CK_TILE_DEVICE auto GetABLdsTensorViews(void* p_smem) const { // A tile in LDS @@ -91,8 +71,8 @@ struct GemmPipelineAgBgCrImplBase return make_tuple(std::move(a_lds_block), std::move(b_lds_block)); } - template - CK_TILE_DEVICE constexpr auto GetAWindows(const ADramBlockWindowTmp& a_dram_block_window_tmp, + template + CK_TILE_DEVICE constexpr auto GetAWindows(const ADramBlockWindow& a_dram_block_window, const ALdsTensorView& a_lds_block_view, const ALdsLoadTileDistr&) const { @@ -105,16 +85,16 @@ struct GemmPipelineAgBgCrImplBase decltype(auto) a_copy_dram_window = [&] { if constexpr(Problem::SkipALds == false) { - return make_tile_window(a_dram_block_window_tmp.get_bottom_tensor_view(), + return make_tile_window(a_dram_block_window.get_bottom_tensor_view(), make_tuple(YPerTile{}, XPerTile{}), - a_dram_block_window_tmp.get_window_origin(), + a_dram_block_window.get_window_origin(), Policy::template MakeADramTileDistribution()); } else { - return make_tile_window(a_dram_block_window_tmp.get_bottom_tensor_view(), + return make_tile_window(a_dram_block_window.get_bottom_tensor_view(), make_tuple(YPerTile{}, XPerTile{}), - a_dram_block_window_tmp.get_window_origin(), + a_dram_block_window.get_window_origin(), ALdsLoadTileDistr{}); } }(); @@ -134,8 +114,8 @@ struct GemmPipelineAgBgCrImplBase std::move(a_lds_gemm_window)); } - template - CK_TILE_DEVICE constexpr auto GetBWindows(const BDramBlockWindowTmp& b_dram_block_window_tmp, + template + CK_TILE_DEVICE constexpr auto GetBWindows(const BDramBlockWindow& b_dram_block_window, const BLdsTensorView& b_lds_block_view, const BLdsLoadTileDistr&) const { @@ -147,16 +127,16 @@ struct GemmPipelineAgBgCrImplBase decltype(auto) b_copy_dram_window = [&] { if constexpr(Problem::SkipBLds == false) { - return make_tile_window(b_dram_block_window_tmp.get_bottom_tensor_view(), + return make_tile_window(b_dram_block_window.get_bottom_tensor_view(), make_tuple(YPerTile{}, XPerTile{}), - b_dram_block_window_tmp.get_window_origin(), + b_dram_block_window.get_window_origin(), Policy::template MakeBDramTileDistribution()); } else { - return make_tile_window(b_dram_block_window_tmp.get_bottom_tensor_view(), + return make_tile_window(b_dram_block_window.get_bottom_tensor_view(), make_tuple(YPerTile{}, XPerTile{}), - b_dram_block_window_tmp.get_window_origin(), + b_dram_block_window.get_window_origin(), BLdsLoadTileDistr{}); } }(); @@ -177,42 +157,6 @@ struct GemmPipelineAgBgCrImplBase std::move(b_copy_lds_window), std::move(b_lds_gemm_window)); } - - template - CK_TILE_DEVICE constexpr auto - GetADramWindowSkipLds(const ADramBlockWindowTmp& a_dram_block_window_tmp, - const ALdsLoadTileDistr&) const - { - constexpr bool is_col_major = std::is_same_v; - - using YPerTile = std::conditional_t, number>; - using XPerTile = std::conditional_t, number>; - - auto a_copy_dram_window = make_tile_window(a_dram_block_window_tmp.get_bottom_tensor_view(), - make_tuple(YPerTile{}, XPerTile{}), - a_dram_block_window_tmp.get_window_origin(), - ALdsLoadTileDistr{}); - - return a_copy_dram_window; - } - - template - CK_TILE_DEVICE constexpr auto - GetBDramWindowSkipLds(const BDramBlockWindowTmp& b_dram_block_window_tmp, - const BLdsLoadTileDistr&) const - { - constexpr bool is_row_major = std::is_same_v; - - using YPerTile = std::conditional_t, number>; - using XPerTile = std::conditional_t, number>; - - auto b_copy_dram_window = make_tile_window(b_dram_block_window_tmp.get_bottom_tensor_view(), - make_tuple(YPerTile{}, XPerTile{}), - b_dram_block_window_tmp.get_window_origin(), - BLdsLoadTileDistr{}); - - return b_copy_dram_window; - } }; } // namespace ck_tile diff --git a/test/ck_tile/gemm/test_gemm_pipeline_util.hpp b/test/ck_tile/gemm/test_gemm_pipeline_util.hpp index e37d20c1ef..89ffeae54c 100644 --- a/test/ck_tile/gemm/test_gemm_pipeline_util.hpp +++ b/test/ck_tile/gemm/test_gemm_pipeline_util.hpp @@ -35,8 +35,6 @@ auto calculate_rtol_atol(const ck_tile::index_t K, enum struct GemmPipelineType { Mem, - MemSkipALds, - MemSkipBLds, CompV3, CompV4 };