ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-13 10:37:42 +00:00
Code Issues Packages Projects Releases Wiki Activity

fix wp gemm bug when permuteN is false

Browse Source
This commit is contained in:
valarLip
2025-09-26 08:11:59 +00:00
parent dcd33a6ecc
commit e4391834ad
2 changed files with 3 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -72,6 +72,7 @@ struct GemmConfigBase
static constexpr ck_tile::index_t Pipeline = CK_TILE_PIPELINE_COMPUTE_V3;
static constexpr ck_tile::index_t NumWaveGroups = 1;
static constexpr bool Preshuffle = false;
static constexpr bool TiledMMAPermuteN = false;
};
template <typename PrecType>
@@ -300,7 +301,7 @@ struct GemmConfigPreshufflePrefill : public GemmConfigBase
static constexpr bool Preshuffle = true;
static constexpr bool DoubleSmemBuffer = true;
static constexpr int N_Repeat = N_Tile / N_Warp_Tile / N_Warp;
static constexpr bool TiledMMAPermuteN = N_Repeat % 2 == 0;
static constexpr bool TiledMMAPermuteN = N_Repeat % 4 == 0;
};
template <typename PrecType>

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

@@ -109,7 +109,7 @@ struct WeightPreshuffleInvoker
GemmConfig::NumWaveGroups,
false,
1,
true>>;
GemmConfig::TiledMMAPermuteN>>;
using Kernel = ck_tile::GemmKernel<TilePartitioner, GemmPipeline, GemmEpilogue>;
auto kargs = Kernel::MakeKernelArgs(args);