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

Fix broken ping pong pipeline functionality in the develop branch.

Browse Source
This commit is contained in:
Sudhir Kylasa
2025-09-16 18:31:35 +00:00
parent 7d7ded62d3
commit e1c87eac08
3 changed files with 4 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
example/ck_tile/03_gemm/gemm_utils.hpp
View File

@@ -252,7 +252,7 @@ struct GemmConfigComputeV5 : public GemmConfigBase
static constexpr bool DoubleSmemBuffer = false;
static constexpr ck_tile::index_t Pipeline = CK_TILE_PIPELINE_COMPUTE_V5;
static constexpr ck_tile::index_t NumWaNumWaveGroups = 2;
static constexpr ck_tile::index_t NumWaveGroups = 2;
};
template <typename PrecType>

1
example/ck_tile/03_gemm/universal_gemm_invoker.hpp
View File

@@ -102,6 +102,7 @@ struct UniversalInvoker
TilePartitioner::NPerBlock,
GemmConfig::M_Warp,
GemmConfig::N_Warp,
GemmConfig::K_Warp,
GemmConfig::M_Warp_Tile,
GemmConfig::N_Warp_Tile,
GemmConfig::K_Warp_Tile,

3
include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp
View File

@@ -40,6 +40,7 @@ template <typename ADataType_,
index_t kN_,
index_t MWave_,
index_t NWave_,
index_t KWave_,
index_t MPerXdl_,
index_t NPerXdl_,
index_t KPerXdl_,
@@ -59,7 +60,7 @@ struct CShuffleEpilogueProblem
using DsLayout = remove_cvref_t<DsLayout_>;
using ELayout = remove_cvref_t<ELayout_>;
using CDElementwise = remove_cvref_t<CDElementwise_>;
static constexpr index_t kBlockSize = MWave_ * NWave_ * get_warp_size();
static constexpr index_t kBlockSize = MWave_ * NWave_ * KWave_ * get_warp_size();
static constexpr index_t kMPerBlock = kM_;
static constexpr index_t kNPerBlock = kN_;
static constexpr index_t MWave = MWave_;