/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/fmha/kernel/fmha_fwd_pagedkv_kernel.hpp Source File#
fmha_fwd_pagedkv_kernel.hpp
Go to the documentation of this file.
97 "r" + _TS_(g0br::at(ck_tile::number<0>{})) + "x" + _TS_(g0br::at(ck_tile::number<1>{})) + "x" + _TS_(g0br::at(ck_tile::number<2>{})) + "_" +
98 "r" + _TS_(g1br::at(ck_tile::number<0>{})) + "x" + _TS_(g1br::at(ck_tile::number<1>{})) + "x" + _TS_(g1br::at(ck_tile::number<2>{})) + "_" +
99 "w" + _TS_(g0wt::at(ck_tile::number<0>{})) + "x" + _TS_(g0wt::at(ck_tile::number<1>{})) + "x" + _TS_(g0wt::at(ck_tile::number<2>{})) + "_" +
100 "w" + _TS_(g1wt::at(ck_tile::number<0>{})) + "x" + _TS_(g1wt::at(ck_tile::number<1>{})) + "x" + _TS_(g1wt::at(ck_tile::number<2>{})) + "_" +
101 (kBlockPerCuInput == -1 ? "" : ("o" + _TS_(kBlockPerCu) + "_")) + _SS_(FmhaPipeline::name) + "_" +
102 "v" + (std::is_same_v<VLayout, ck_tile::tensor_layout::gemm::RowMajor> ? "r" : "c") + (pn.empty() ? "_npad" : "_" + pn) +
103 (kHasLogitsSoftCap ? "_logits" : "_nlogits" ) + (BiasEnum == BlockAttentionBiasEnum::NO_BIAS ? _SS_("_nbias") : (_SS_("_") + BlockAttentionBiasEnumToStr<BiasEnum>::name)) +
104 (kHasMask ? "_" + _SS_(FmhaMask::name) : "_nmask") + (kStoreLSE ? "_lse" : "_nlse" ) + (kSkipMinSeqlenQ ? "_skip" : "_nskip" ) + (kDoFp8StaticQuant ? "_squant" : "_nsquant" ) + (kIsPagedKV ? "_pagedkv" : "_npagedkv" );
#define _TS_
#define _SS_
Definition: cluster_descriptor.hpp:13
constexpr CK_TILE_DEVICE auto make_null_tile_window(const WindowLengths &window_lengths)
Definition: null_tile_window.hpp:66
constexpr CK_TILE_HOST_DEVICE auto integer_divide_ceil(X x, Y y)
Definition: math.hpp:149
@ ELEMENTWISE_BIAS
CK_TILE_HOST_DEVICE auto make_page_block_navigator(const TensorView &tensor_view)
Definition: page_block_navigator.hpp:333
constexpr CK_TILE_HOST_DEVICE auto transform_tensor_view(const OldTensorView &old_tensor_view, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition: tensor_view.hpp:511
constexpr CK_TILE_HOST_DEVICE auto pad_tensor_view(const TensorView &tensor_view, const TileLengths &tile_lengths, DoPads)
Definition: tensor_view.hpp:530
constexpr CK_TILE_HOST_DEVICE auto make_pass_through_transform(const LowLength &low_length)
Definition: coordinate_transform.hpp:1558
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
constexpr CK_TILE_DEVICE auto make_tile_window(null_tensor_view, const WindowLengths &window_lengths, const multi_index< WindowLengths::size()> &, Ts &&...)
Definition: null_tile_window.hpp:75
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:360
@ MASK_FROM_TOP_LEFT
@ FROM_BOTTOM_RIGHT
__host__ __device__ composes(Ts &&...) -> composes< remove_cvref_t< Ts >... >
FIXME: create macro to replace 'host device' and nothing more.
typename conditional< predicate, X, Y >::type conditional_t
Definition: functional.hpp:115
Definition: block_position_encoding.hpp:48
Definition: block_attention_bias_enum.hpp:19
Definition: block_position_encoding.hpp:137
Definition: fmha_fwd_pagedkv_kernel.hpp:281
ck_tile::index_t qo_head_idx
Definition: fmha_fwd_pagedkv_kernel.hpp:283
ck_tile::index_t batch_idx
Definition: fmha_fwd_pagedkv_kernel.hpp:282
ck_tile::index_t kv_head_idx
Definition: fmha_fwd_pagedkv_kernel.hpp:284
Definition: fmha_fwd_pagedkv_kernel.hpp:227
const int32_t * cache_batch_idx
Definition: fmha_fwd_pagedkv_kernel.hpp:228
Definition: fmha_fwd_pagedkv_kernel.hpp:215
ck_tile::index_t batch_stride_block_table
Definition: fmha_fwd_pagedkv_kernel.hpp:217
const int32_t * block_table_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:216
ck_tile::index_t page_block_size
Definition: fmha_fwd_pagedkv_kernel.hpp:218
Definition: fmha_fwd_pagedkv_kernel.hpp:183
ck_tile::index_t alibi_slope_stride
Definition: fmha_fwd_pagedkv_kernel.hpp:186
const void * alibi_slope_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:185
Definition: fmha_fwd_pagedkv_kernel.hpp:178
ck_tile::index_t batch_stride_bias
Definition: fmha_fwd_pagedkv_kernel.hpp:179
Definition: fmha_fwd_pagedkv_kernel.hpp:243
ck_tile::index_t batch_stride_k
Definition: fmha_fwd_pagedkv_kernel.hpp:247
ck_tile::index_t batch_stride_q
Definition: fmha_fwd_pagedkv_kernel.hpp:246
const int32_t * seqlen_k_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:244
ck_tile::index_t batch_stride_o
Definition: fmha_fwd_pagedkv_kernel.hpp:251
ck_tile::index_t batch_stride_v
Definition: fmha_fwd_pagedkv_kernel.hpp:249
Definition: fmha_fwd_pagedkv_kernel.hpp:171
ck_tile::index_t stride_bias
Definition: fmha_fwd_pagedkv_kernel.hpp:173
const void * bias_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:172
ck_tile::index_t nhead_stride_bias
Definition: fmha_fwd_pagedkv_kernel.hpp:174
Definition: fmha_fwd_pagedkv_kernel.hpp:120
ck_tile::index_t hdim_v
Definition: fmha_fwd_pagedkv_kernel.hpp:129
ck_tile::index_t seqlen_q
Definition: fmha_fwd_pagedkv_kernel.hpp:126
ck_tile::index_t nhead_stride_k
Definition: fmha_fwd_pagedkv_kernel.hpp:143
ck_tile::index_t stride_o
Definition: fmha_fwd_pagedkv_kernel.hpp:140
const void * k_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:122
ck_tile::index_t nhead_stride_q
Definition: fmha_fwd_pagedkv_kernel.hpp:142
ck_tile::index_t stride_v
Definition: fmha_fwd_pagedkv_kernel.hpp:139
float scale_s
Definition: fmha_fwd_pagedkv_kernel.hpp:135
ck_tile::index_t stride_k
Definition: fmha_fwd_pagedkv_kernel.hpp:138
const void * v_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:123
ck_tile::index_t nhead_stride_v
Definition: fmha_fwd_pagedkv_kernel.hpp:144
ck_tile::index_t nhead_stride_o
Definition: fmha_fwd_pagedkv_kernel.hpp:145
ck_tile::index_t hdim_q
Definition: fmha_fwd_pagedkv_kernel.hpp:128
ck_tile::index_t seqlen_k
Definition: fmha_fwd_pagedkv_kernel.hpp:127
const void * q_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:121
ck_tile::index_t num_head_q
Definition: fmha_fwd_pagedkv_kernel.hpp:131
ck_tile::index_t nhead_ratio_qk
Definition: fmha_fwd_pagedkv_kernel.hpp:134
void * o_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:124
ck_tile::index_t stride_q
Definition: fmha_fwd_pagedkv_kernel.hpp:137
Definition: fmha_fwd_pagedkv_kernel.hpp:203
ck_tile::index_t batch_stride_lse
Definition: fmha_fwd_pagedkv_kernel.hpp:206
void * lse_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:204
ck_tile::index_t nhead_stride_lse
Definition: fmha_fwd_pagedkv_kernel.hpp:205
Definition: fmha_fwd_pagedkv_kernel.hpp:113
Definition: fmha_fwd_pagedkv_kernel.hpp:197
float scale_p
Definition: fmha_fwd_pagedkv_kernel.hpp:198
float scale_o
Definition: fmha_fwd_pagedkv_kernel.hpp:199
Definition: fmha_fwd_pagedkv_kernel.hpp:267
const int32_t * seqlen_k_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:270
const int32_t * seqstart_k_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:269
ck_tile::index_t batch_stride_k
Definition: fmha_fwd_pagedkv_kernel.hpp:272
ck_tile::index_t batch_stride_v
Definition: fmha_fwd_pagedkv_kernel.hpp:274
const int32_t * seqstart_q_ptr
Definition: fmha_fwd_pagedkv_kernel.hpp:268
Definition: fmha_fwd_pagedkv_kernel.hpp:149
FmhaFwdLogitsSoftCapKargs()=default
float logits_soft_cap
Definition: fmha_fwd_pagedkv_kernel.hpp:166
void init_logits_soft_cap(float logits_soft_cap_)
Definition: fmha_fwd_pagedkv_kernel.hpp:152
float logits_soft_cap_rcp
Definition: fmha_fwd_pagedkv_kernel.hpp:167
Definition: fmha_fwd_pagedkv_kernel.hpp:190
ck_tile::index_t window_size_left
Definition: fmha_fwd_pagedkv_kernel.hpp:192
ck_tile::index_t window_size_right
Definition: fmha_fwd_pagedkv_kernel.hpp:192
ck_tile::GenericAttentionMaskEnum mask_type
Definition: fmha_fwd_pagedkv_kernel.hpp:193
Definition: fmha_fwd_pagedkv_kernel.hpp:210
ck_tile::index_t min_seqlen_q
Definition: fmha_fwd_pagedkv_kernel.hpp:211
Definition: fmha_fwd_pagedkv_kernel.hpp:222
bool is_gappy
Definition: fmha_fwd_pagedkv_kernel.hpp:223
Definition: fmha_fwd_pagedkv_kernel.hpp:66
Definition: fmha_fwd_pagedkv_kernel.hpp:28
static constexpr CK_TILE_HOST std::enable_if_t< Cond, Kargs > MakeKargsImpl(const void *q_ptr, const void *k_ptr, const void *v_ptr, const void *bias_ptr, void *lse_ptr, void *o_ptr, ck_tile::index_t seqlen_q, ck_tile::index_t seqlen_k, const void *seqlen_k_ptr, ck_tile::index_t hdim_q, ck_tile::index_t hdim_v, ck_tile::index_t num_head_q, ck_tile::index_t nhead_ratio_qk, const void *block_table_ptr, ck_tile::index_t batch_stride_block_table, ck_tile::index_t page_block_size, const void *cache_batch_idx, float scale_s, float scale_p, float scale_o, float logits_soft_cap, ck_tile::index_t stride_q, ck_tile::index_t stride_k, ck_tile::index_t stride_v, ck_tile::index_t stride_bias, ck_tile::index_t stride_o, ck_tile::index_t nhead_stride_q, ck_tile::index_t nhead_stride_k, ck_tile::index_t nhead_stride_v, ck_tile::index_t nhead_stride_bias, ck_tile::index_t nhead_stride_lse, ck_tile::index_t nhead_stride_o, ck_tile::index_t batch_stride_q, ck_tile::index_t batch_stride_k, ck_tile::index_t batch_stride_v, ck_tile::index_t batch_stride_bias, ck_tile::index_t batch_stride_lse, ck_tile::index_t batch_stride_o, ck_tile::index_t window_size_left, ck_tile::index_t window_size_right, ck_tile::index_t mask_type)
Definition: fmha_fwd_pagedkv_kernel.hpp:289
static constexpr bool kIsGroupMode
Definition: fmha_fwd_pagedkv_kernel.hpp:47
static constexpr bool kHasLogitsSoftCap
Definition: fmha_fwd_pagedkv_kernel.hpp:52
ck_tile::remove_cvref_t< EpiloguePipeline_ > EpiloguePipeline
Definition: fmha_fwd_pagedkv_kernel.hpp:30
static constexpr ck_tile::index_t kBlockPerCu
Definition: fmha_fwd_pagedkv_kernel.hpp:32
static constexpr bool kStoreLSE
Definition: fmha_fwd_pagedkv_kernel.hpp:54
static CK_TILE_HOST std::string GetName()
Definition: fmha_fwd_pagedkv_kernel.hpp:74
static constexpr bool kPadSeqLenK
Definition: fmha_fwd_pagedkv_kernel.hpp:49
static constexpr bool kIsPagedKV
Definition: fmha_fwd_pagedkv_kernel.hpp:57
static constexpr CK_TILE_HOST std::enable_if_t< Cond, Kargs > MakeKargs(const void *q_ptr, const void *k_ptr, const void *v_ptr, const void *bias_ptr, void *lse_ptr, void *o_ptr, const void *seqstart_q_ptr, const void *seqstart_k_ptr, const void *seqlen_k_ptr, ck_tile::index_t hdim_q, ck_tile::index_t hdim_v, ck_tile::index_t num_head_q, ck_tile::index_t nhead_ratio_qk, const void *block_table_ptr, ck_tile::index_t batch_stride_block_table, ck_tile::index_t page_block_size, bool is_gappy, float scale_s, float scale_p, float scale_o, float logits_soft_cap, ck_tile::index_t stride_q, ck_tile::index_t stride_k, ck_tile::index_t stride_v, ck_tile::index_t stride_bias, ck_tile::index_t stride_o, ck_tile::index_t nhead_stride_q, ck_tile::index_t nhead_stride_k, ck_tile::index_t nhead_stride_v, ck_tile::index_t nhead_stride_bias, ck_tile::index_t nhead_stride_lse, ck_tile::index_t nhead_stride_o, ck_tile::index_t batch_stride_k, ck_tile::index_t batch_stride_v, ck_tile::index_t window_size_left, ck_tile::index_t window_size_right, ck_tile::index_t mask_type, ck_tile::index_t min_seqlen_q)
Definition: fmha_fwd_pagedkv_kernel.hpp:627
static constexpr bool kSkipMinSeqlenQ
Definition: fmha_fwd_pagedkv_kernel.hpp:56
static CK_TILE_HOST void PrintParameters(const Kargs &kargs, int num_batches)
Definition: fmha_fwd_pagedkv_kernel.hpp:706
static constexpr CK_TILE_HOST std::enable_if_t< Cond, Kargs > MakeKargs(const void *q_ptr, const void *k_ptr, const void *v_ptr, const void *bias_ptr, void *lse_ptr, void *o_ptr, ck_tile::index_t seqlen_q, ck_tile::index_t seqlen_k, const void *seqlen_k_ptr, ck_tile::index_t hdim_q, ck_tile::index_t hdim_v, ck_tile::index_t num_head_q, ck_tile::index_t nhead_ratio_qk, const void *block_table_ptr, ck_tile::index_t batch_stride_block_table, ck_tile::index_t page_block_size, const void *cache_batch_idx, float scale_s, float scale_p, float scale_o, float logits_soft_cap, ck_tile::index_t stride_q, ck_tile::index_t stride_k, ck_tile::index_t stride_v, ck_tile::index_t stride_bias, ck_tile::index_t stride_o, ck_tile::index_t nhead_stride_q, ck_tile::index_t nhead_stride_k, ck_tile::index_t nhead_stride_v, ck_tile::index_t nhead_stride_bias, ck_tile::index_t nhead_stride_lse, ck_tile::index_t nhead_stride_o, ck_tile::index_t batch_stride_q, ck_tile::index_t batch_stride_k, ck_tile::index_t batch_stride_v, ck_tile::index_t batch_stride_bias, ck_tile::index_t batch_stride_lse, ck_tile::index_t batch_stride_o, ck_tile::index_t window_size_left, ck_tile::index_t window_size_right, ck_tile::index_t mask_type)
Definition: fmha_fwd_pagedkv_kernel.hpp:416
static constexpr ck_tile::index_t kBlockSize
Definition: fmha_fwd_pagedkv_kernel.hpp:31
static constexpr bool kPadHeadDimV
Definition: fmha_fwd_pagedkv_kernel.hpp:51
ck_tile::remove_cvref_t< typename FmhaPipeline::LSEDataType > LSEDataType
Definition: fmha_fwd_pagedkv_kernel.hpp:41
ck_tile::remove_cvref_t< typename FmhaPipeline::VDataType > VDataType
Definition: fmha_fwd_pagedkv_kernel.hpp:39
static constexpr CK_TILE_DEVICE auto GetTileIndex(const Kargs &kargs)
Definition: fmha_fwd_pagedkv_kernel.hpp:803
static constexpr ck_tile::index_t kBlockPerCuInput
Definition: fmha_fwd_pagedkv_kernel.hpp:35
static constexpr bool kUseAsyncCopy
Definition: fmha_fwd_pagedkv_kernel.hpp:63
static constexpr bool kHasMask
Definition: fmha_fwd_pagedkv_kernel.hpp:61
static constexpr CK_TILE_HOST auto GridSize(ck_tile::index_t batch_size_, ck_tile::index_t nhead_, ck_tile::index_t seqlen_q_, ck_tile::index_t hdim_v_, bool has_padded_seqlen_k)
Definition: fmha_fwd_pagedkv_kernel.hpp:778
ck_tile::remove_cvref_t< FmhaPipeline_ > FmhaPipeline
Definition: fmha_fwd_pagedkv_kernel.hpp:29
ck_tile::remove_cvref_t< typename FmhaPipeline::ODataType > ODataType
Definition: fmha_fwd_pagedkv_kernel.hpp:42
static constexpr CK_TILE_HOST auto BlockSize()
Definition: fmha_fwd_pagedkv_kernel.hpp:868
ck_tile::remove_cvref_t< typename FmhaPipeline::BiasDataType > BiasDataType
Definition: fmha_fwd_pagedkv_kernel.hpp:40
static constexpr bool kPadHeadDimQ
Definition: fmha_fwd_pagedkv_kernel.hpp:50
ck_tile::remove_cvref_t< typename FmhaPipeline::VLayout > VLayout
Definition: fmha_fwd_pagedkv_kernel.hpp:45
static constexpr bool kPadSeqLenQ
Definition: fmha_fwd_pagedkv_kernel.hpp:48
ck_tile::remove_cvref_t< typename FmhaPipeline::SaccDataType > SaccDataType
Definition: fmha_fwd_pagedkv_kernel.hpp:43
ck_tile::remove_cvref_t< typename FmhaPipeline::FmhaMask > FmhaMask
Definition: fmha_fwd_pagedkv_kernel.hpp:60
static constexpr CK_TILE_HOST_DEVICE ck_tile::index_t GetSmemSize()
Definition: fmha_fwd_pagedkv_kernel.hpp:870
CK_TILE_DEVICE void operator()(Kargs kargs) const
Definition: fmha_fwd_pagedkv_kernel.hpp:875
ck_tile::remove_cvref_t< typename FmhaPipeline::AttentionVariant > AttentionVariant
Definition: fmha_fwd_pagedkv_kernel.hpp:59
ck_tile::remove_cvref_t< typename FmhaPipeline::KDataType > KDataType
Definition: fmha_fwd_pagedkv_kernel.hpp:38
ck_tile::remove_cvref_t< typename FmhaPipeline::QDataType > QDataType
Definition: fmha_fwd_pagedkv_kernel.hpp:37
static constexpr bool kDoFp8StaticQuant
Definition: fmha_fwd_pagedkv_kernel.hpp:55
std::conditional_t< kIsGroupMode, FmhaFwdGroupModeKargs, FmhaFwdBatchModeKargs > Kargs
Definition: fmha_fwd_pagedkv_kernel.hpp:278
static constexpr CK_TILE_HOST std::enable_if_t< Cond, Kargs > MakeKargsImpl(const void *q_ptr, const void *k_ptr, const void *v_ptr, const void *bias_ptr, void *lse_ptr, void *o_ptr, const void *seqstart_q_ptr, const void *seqstart_k_ptr, const void *seqlen_k_ptr, ck_tile::index_t hdim_q, ck_tile::index_t hdim_v, ck_tile::index_t num_head_q, ck_tile::index_t nhead_ratio_qk, const void *block_table_ptr, ck_tile::index_t batch_stride_block_table, ck_tile::index_t page_block_size, bool is_gappy, float scale_s, float scale_p, float scale_o, float logits_soft_cap, ck_tile::index_t stride_q, ck_tile::index_t stride_k, ck_tile::index_t stride_v, ck_tile::index_t stride_bias, ck_tile::index_t stride_o, ck_tile::index_t nhead_stride_q, ck_tile::index_t nhead_stride_k, ck_tile::index_t nhead_stride_v, ck_tile::index_t nhead_stride_bias, ck_tile::index_t nhead_stride_lse, ck_tile::index_t nhead_stride_o, ck_tile::index_t batch_stride_k, ck_tile::index_t batch_stride_v, ck_tile::index_t window_size_left, ck_tile::index_t window_size_right, ck_tile::index_t mask_type, ck_tile::index_t min_seqlen_q)
Definition: fmha_fwd_pagedkv_kernel.hpp:503
static constexpr auto BiasEnum
Definition: fmha_fwd_pagedkv_kernel.hpp:53
Definition: variants.hpp:63
Definition: variants.hpp:51
Definition: integral_constant.hpp:13
Definition: functional.hpp:86
Definition: unary_element_function.hpp:56
Definition: math.hpp:28
Definition: sequence.hpp:49