diff --git a/include/ck_tile/ops/fmha/kernel/fmha_fwd_pagedkv_kernel.hpp b/include/ck_tile/ops/fmha/kernel/fmha_fwd_pagedkv_kernel.hpp index 01ce70fffb..7ddfbb733c 100644 --- a/include/ck_tile/ops/fmha/kernel/fmha_fwd_pagedkv_kernel.hpp +++ b/include/ck_tile/ops/fmha/kernel/fmha_fwd_pagedkv_kernel.hpp @@ -1016,14 +1016,14 @@ struct FmhaFwdPagedKVKernel return pad_tensor_view( q_dram_naive, make_tuple(number{}, number{}), - sequence{}); + sequence{}); } else { return pad_tensor_view( q_dram_naive, make_tuple(number{}, number{}), - sequence{}); + sequence{}); } }(); @@ -1120,7 +1120,7 @@ struct FmhaFwdPagedKVKernel const index_t num_blocks = integer_divide_ceil(kv_l2p_offset + kargs.seqlen_k, kargs.page_block_size); - const long_index_t fixed_offset = i_nhead_ * kargs.nhead_stride_k; + const long_index_t fixed_offset = static_cast(i_nhead_) * kargs.nhead_stride_k; return make_page_block_navigator( kargs.k_ptr, @@ -1150,7 +1150,7 @@ struct FmhaFwdPagedKVKernel const index_t num_blocks = integer_divide_ceil(kv_l2p_offset + kargs.seqlen_k, kargs.page_block_size); - const long_index_t fixed_offset = i_nhead_ * kargs.nhead_stride_v; + const long_index_t fixed_offset = static_cast(i_nhead_) * kargs.nhead_stride_v; return make_page_block_navigator( kargs.v_ptr, @@ -1197,7 +1197,7 @@ struct FmhaFwdPagedKVKernel return pad_tensor_view(bias_dram_naive, bias_dram_window_lengths, - sequence{}); + sequence{}); }(); return make_tile_window(bias_dram, bias_dram_window_lengths, {i_m0, 0}); diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_problem.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_problem.hpp index 72060c2272..f868804cc2 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_problem.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_problem.hpp @@ -110,7 +110,6 @@ struct BlockFmhaFwdPagedKVPipelineProblem static constexpr bool kStoreLSE = Traits::kStoreLSE; static constexpr bool kDoFp8StaticQuant = Traits::kDoFp8StaticQuant; static constexpr bool kIsPagedKV = Traits::kIsPagedKV; - static constexpr bool kMergeNumHeadGroupsSeqLenQ = Traits::kMergeNumHeadGroupsSeqLenQ; static constexpr index_t kBlockPerCu = Traits::kBlockPerCu; };