ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-05 20:55:59 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add constraint to kMaxSplits

Browse Source
This commit is contained in:
PoYen, Chen
2024-06-12 05:00:06 +00:00
parent e00ff9d246
commit 0a2132d758
2 changed files with 4 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_combine_pipeline_default_policy.hpp
View File

@@ -137,7 +137,8 @@ struct BlockFmhaFwdSplitKVCombinePipelineDefaultPolicy
{
constexpr index_t kBlockSize = Problem::kBlockSize;
constexpr index_t kNPerBlock = max(Problem::kMaxSplits, get_warp_size());
static_assert(0 < Problem::kMaxSplits && Problem::kMaxSplits % get_warp_size() == 0);
constexpr index_t kNPerBlock = Problem::kMaxSplits;
constexpr index_t kMPerBlock = Problem::kM0;
constexpr index_t NumElements = (kMPerBlock * kNPerBlock);

3
include/ck_tile/ops/fmha/pipeline/tile_fmha_traits.hpp
View File

@@ -45,7 +45,8 @@ struct TileFmhaFwdSplitKVCombineTraits
static constexpr bool kStoreLSE = kStoreLSE_;
static constexpr bool kDoFp8StaticQuant = kDoFp8StaticQuant_;
static constexpr index_t kMaxSplits = kMaxSplits_;
static constexpr index_t kMaxSplits = kMaxSplits_;
static_assert(0 < kMaxSplits && kMaxSplits % get_warp_size() == 0);
static constexpr index_t kBlockPerCu = kBlockPerCu_;
};