ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-03-23 08:37:39 +00:00
Code Issues Packages Projects Releases Wiki Activity

optimize gemm2 atomic_add pattern

Browse Source
This commit is contained in:
Feng Shijie
2025-08-11 08:38:23 +00:00
parent 87aed564dc
commit cc9c7b9e58
1 changed files with 8 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
include/ck_tile/ops/moe_flatmm/kernel/moe_flatmm_kernel.hpp
View File

@@ -693,13 +693,14 @@ struct MoeFlatmmKernel
"Currently, the CShuffle EpiloguePipeline only supports the Row Major "
"Output layout");
using TileEncodingPattern =
TileDistributionEncodingPattern2D<kBlockSize,
MPerIterationShuffle,
NPerIterationShuffle,
EpiloguePipeline::GetVectorSizeC(),
tile_distribution_pattern::thread_raked,
EpiProblem::kNumWaveGroups>;
using TileEncodingPattern = TileDistributionEncodingPattern2D<
kBlockSize,
MPerIterationShuffle,
NPerIterationShuffle,
kind == MoeFlatmmKind::kFFN_gemm2 ? 2 : EpiloguePipeline::GetVectorSizeC(),
tile_distribution_pattern::thread_raked,
EpiProblem::kNumWaveGroups>;
constexpr auto dram_tile_distribution =
TileEncodingPattern::Make2DStaticTileDistribution();