BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ > Struct Template Reference

BlockFmhaFwdAppendKVPipeline&lt; Problem_, Policy_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Attributes | List of all members
ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ > Struct Template Reference

#include <block_fmha_fwd_appendkv_pipeline.hpp>

Public Types

using Problem = remove_cvref_t<Problem_>
using Policy = remove_cvref_t<Policy_>
using QDataType = typename Problem::QDataType
using KDataType = typename Problem::KDataType
using VDataType = typename Problem::VDataType
using VLayout = typename Problem::VLayout

Public Member Functions

template<typename QDramBlockWindow, typename KDramBlockWindow, typename KPageBlockNavigator, typename KnewDramBlockWindow, typename VDramBlockWindow, typename VPageBlockNavigator, typename VnewDramBlockWindow, typename QElementFunction, typename KnewElementFunction, typename VnewElementFunction, typename QRotaryCosDramBlockWindow, typename QRotarySinDramBlockWindow, typename KnewRotaryCosDramBlockWindow, typename KnewRotarySinDramBlockWindow>
CK_TILE_HOST_DEVICE auto operator() (QDramBlockWindow &q_dram_block_window, const QElementFunction &q_element_func, KDramBlockWindow &k_dram_block_window, index_t i_page_block_k, const KPageBlockNavigator &k_page_block_navigator, const KnewDramBlockWindow &knew_dram_block_window, const KnewElementFunction &knew_element_func, VDramBlockWindow &v_dram_block_window, index_t i_page_block_v, const VPageBlockNavigator &v_page_block_navigator, const VnewDramBlockWindow &vnew_dram_block_window, const VnewElementFunction &vnew_element_func, const QRotaryCosDramBlockWindow q_rotary_cos_dram_block_window, const QRotarySinDramBlockWindow q_rotary_sin_dram_block_window, const KnewRotaryCosDramBlockWindow knew_rotary_cos_dram_block_window, const KnewRotarySinDramBlockWindow knew_rotary_sin_dram_block_window, index_t rotary_dim, bool skip_rotate_q, bool skip_rotate_append_kv) const
template<typename QDramBlockWindow, typename KDramBlockWindow, typename KPageBlockNavigator, typename KnewDramBlockWindow, typename VDramBlockWindow, typename VPageBlockNavigator, typename VnewDramBlockWindow, typename QRotaryCosDramBlockWindow, typename QRotarySinDramBlockWindow, typename KnewRotaryCosDramBlockWindow, typename KnewRotarySinDramBlockWindow>
CK_TILE_HOST_DEVICE auto operator() (QDramBlockWindow &q_dram_block_window, KDramBlockWindow &k_dram_block_window, index_t i_page_block_k, const KPageBlockNavigator &k_page_block_navigator, const KnewDramBlockWindow &knew_dram_block_window, VDramBlockWindow &v_dram_block_window, index_t i_page_block_v, const VPageBlockNavigator &v_page_block_navigator, const VnewDramBlockWindow &vnew_dram_block_window, const QRotaryCosDramBlockWindow &q_rotary_cos_dram_block_window, const QRotarySinDramBlockWindow &q_rotary_sin_dram_block_window, const KnewRotaryCosDramBlockWindow &knew_rotary_cos_dram_block_window, const KnewRotarySinDramBlockWindow &knew_rotary_sin_dram_block_window, index_t rotary_dim, bool skip_rotate_q, bool skip_rotate_append_kv) const

Static Public Attributes

static constexpr index_t kBlockSize = Problem::kBlockSize
static constexpr index_t kM0 = Problem::kM0
static constexpr index_t kN0 = Problem::kN0
static constexpr index_t kK0 = Problem::kK0
static constexpr index_t kN1 = Problem::kN1
static constexpr auto RotaryEnum = Problem::RotaryEnum
static constexpr bool kIsPagedKV = Problem::kIsPagedKV
static constexpr bool kPadSeqLenQ = Problem::kPadSeqLenQ
static constexpr bool kPadSeqLenK = Problem::kPadSeqLenK
static constexpr bool kPadHeadDimQ = Problem::kPadHeadDimQ
static constexpr bool kPadHeadDimV = Problem::kPadHeadDimV
static constexpr index_t kAlignmentQ
static constexpr index_t kAlignmentK
static constexpr index_t kAlignmentV
static constexpr index_t kBlockPerCu

Member Typedef Documentation

◆ KDataType

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
using ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::KDataType = typename Problem::KDataType

◆ Policy

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
using ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::Policy = remove_cvref_t<Policy_>

◆ Problem

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
using ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::Problem = remove_cvref_t<Problem_>

◆ QDataType

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
using ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::QDataType = typename Problem::QDataType

◆ VDataType

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
using ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::VDataType = typename Problem::VDataType

◆ VLayout

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
using ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::VLayout = typename Problem::VLayout

Member Function Documentation

◆ operator()() [1/2]

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
template<typename QDramBlockWindow, typename KDramBlockWindow, typename KPageBlockNavigator, typename KnewDramBlockWindow, typename VDramBlockWindow, typename VPageBlockNavigator, typename VnewDramBlockWindow, typename QElementFunction, typename KnewElementFunction, typename VnewElementFunction, typename QRotaryCosDramBlockWindow, typename QRotarySinDramBlockWindow, typename KnewRotaryCosDramBlockWindow, typename KnewRotarySinDramBlockWindow>
CK_TILE_HOST_DEVICE auto ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::operator() ( QDramBlockWindow & q_dram_block_window,
const QElementFunction & q_element_func,
KDramBlockWindow & k_dram_block_window,
index_t i_page_block_k,
const KPageBlockNavigator & k_page_block_navigator,
const KnewDramBlockWindow & knew_dram_block_window,
const KnewElementFunction & knew_element_func,
VDramBlockWindow & v_dram_block_window,
index_t i_page_block_v,
const VPageBlockNavigator & v_page_block_navigator,
const VnewDramBlockWindow & vnew_dram_block_window,
const VnewElementFunction & vnew_element_func,
const QRotaryCosDramBlockWindow q_rotary_cos_dram_block_window,
const QRotarySinDramBlockWindow q_rotary_sin_dram_block_window,
const KnewRotaryCosDramBlockWindow knew_rotary_cos_dram_block_window,
const KnewRotarySinDramBlockWindow knew_rotary_sin_dram_block_window,
index_t rotary_dim,
bool skip_rotate_q,
bool skip_rotate_append_kv ) const
inline

◆ operator()() [2/2]

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
template<typename QDramBlockWindow, typename KDramBlockWindow, typename KPageBlockNavigator, typename KnewDramBlockWindow, typename VDramBlockWindow, typename VPageBlockNavigator, typename VnewDramBlockWindow, typename QRotaryCosDramBlockWindow, typename QRotarySinDramBlockWindow, typename KnewRotaryCosDramBlockWindow, typename KnewRotarySinDramBlockWindow>
CK_TILE_HOST_DEVICE auto ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::operator() ( QDramBlockWindow & q_dram_block_window,
KDramBlockWindow & k_dram_block_window,
index_t i_page_block_k,
const KPageBlockNavigator & k_page_block_navigator,
const KnewDramBlockWindow & knew_dram_block_window,
VDramBlockWindow & v_dram_block_window,
index_t i_page_block_v,
const VPageBlockNavigator & v_page_block_navigator,
const VnewDramBlockWindow & vnew_dram_block_window,
const QRotaryCosDramBlockWindow & q_rotary_cos_dram_block_window,
const QRotarySinDramBlockWindow & q_rotary_sin_dram_block_window,
const KnewRotaryCosDramBlockWindow & knew_rotary_cos_dram_block_window,
const KnewRotarySinDramBlockWindow & knew_rotary_sin_dram_block_window,
index_t rotary_dim,
bool skip_rotate_q,
bool skip_rotate_append_kv ) const
inline

Member Data Documentation

◆ kAlignmentK

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kAlignmentK
staticconstexpr
Initial value:
=
kPadHeadDimQ ? 1 : Policy::template GetAlignmentK<Problem>()
ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVR::kPadHeadDimQ
static constexpr index_t kPadHeadDimQ
Definition block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr.hpp:52

◆ kAlignmentQ

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kAlignmentQ
staticconstexpr
Initial value:
=
kPadHeadDimQ ? 1 : Policy::template GetAlignmentQ<Problem>()

◆ kAlignmentV

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kAlignmentV
staticconstexpr
Initial value:
= []() {
if constexpr(std::is_same_v<VLayout, ck_tile::tensor_layout::gemm::RowMajor>)
return kPadHeadDimV ? 1 : Policy::template GetAlignmentV<Problem>();
else
return kPadSeqLenK ? 1 : Policy::template GetAlignmentV<Problem>();
}()
ck_tile::BlockFmhaBatchPrefillPipelineQRKSVSAsync::kPadSeqLenK
static constexpr bool kPadSeqLenK
Definition block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp:64
ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVR::kPadHeadDimV
static constexpr index_t kPadHeadDimV
Definition block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr.hpp:53
ck_tile::BlockFmhaFwdAppendKVPipeline::Policy
remove_cvref_t< Policy_ > Policy
Definition block_fmha_fwd_appendkv_pipeline.hpp:16
ck_tile::BlockFmhaFwdAppendKVPipeline::Problem
remove_cvref_t< Problem_ > Problem
Definition block_fmha_fwd_appendkv_pipeline.hpp:15

◆ kBlockPerCu

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kBlockPerCu
staticconstexpr
Initial value:
= []() {
if constexpr(Problem::kBlockPerCu != -1)
return Problem::kBlockPerCu;
else
{
if constexpr(kK0 <= 32)
{
return 2;
}
else if constexpr(kK0 <= 64)
{
return 3;
}
else if constexpr(kK0 <= 128)
{
return 2;
}
else if constexpr(kK0 <= 256)
{
return 1;
}
}
}()
ck_tile::BlockFmhaFwdAppendKVPipeline::kK0
static constexpr index_t kK0
Definition block_fmha_fwd_appendkv_pipeline.hpp:27

◆ kBlockSize

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kBlockSize = Problem::kBlockSize
staticconstexpr

◆ kIsPagedKV

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
bool ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kIsPagedKV = Problem::kIsPagedKV
staticconstexpr

◆ kK0

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kK0 = Problem::kK0
staticconstexpr

◆ kM0

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kM0 = Problem::kM0
staticconstexpr

◆ kN0

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kN0 = Problem::kN0
staticconstexpr

◆ kN1

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kN1 = Problem::kN1
staticconstexpr

◆ kPadHeadDimQ

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
bool ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kPadHeadDimQ = Problem::kPadHeadDimQ
staticconstexpr

◆ kPadHeadDimV

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
bool ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kPadHeadDimV = Problem::kPadHeadDimV
staticconstexpr

◆ kPadSeqLenK

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
bool ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kPadSeqLenK = Problem::kPadSeqLenK
staticconstexpr

◆ kPadSeqLenQ

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
bool ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::kPadSeqLenQ = Problem::kPadSeqLenQ
staticconstexpr

◆ RotaryEnum

template<typename Problem_, typename Policy_ = BlockFmhaFwdAppendKVPipelineDefaultPolicy>
auto ck_tile::BlockFmhaFwdAppendKVPipeline< Problem_, Policy_ >::RotaryEnum = Problem::RotaryEnum
staticconstexpr
The documentation for this struct was generated from the following file: