/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/fmha/kernel/fmha_batch_prefill_kernel.hpp Source File#
fmha_batch_prefill_kernel.hpp
Go to the documentation of this file.
95 "r" + _TS_(g0br::at(ck_tile::number<0>{})) + "x" + _TS_(g0br::at(ck_tile::number<1>{})) + "x" + _TS_(g0br::at(ck_tile::number<2>{})) + "_" +
96 "r" + _TS_(g1br::at(ck_tile::number<0>{})) + "x" + _TS_(g1br::at(ck_tile::number<1>{})) + "x" + _TS_(g1br::at(ck_tile::number<2>{})) + "_" +
97 "w" + _TS_(g0wt::at(ck_tile::number<0>{})) + "x" + _TS_(g0wt::at(ck_tile::number<1>{})) + "x" + _TS_(g0wt::at(ck_tile::number<2>{})) + "_" +
98 "w" + _TS_(g1wt::at(ck_tile::number<0>{})) + "x" + _TS_(g1wt::at(ck_tile::number<1>{})) + "x" + _TS_(g1wt::at(ck_tile::number<2>{})) + "_" +
99 (kBlockPerCuInput == -1 ? "" : ("o" + _TS_(kBlockPerCu) + "_")) + _SS_(FmhaPipeline::name) + "_" +
100 "v" + (std::is_same_v<VLayout, ck_tile::tensor_layout::gemm::RowMajor> ? "r" : "c") + (pn.empty() ? "_npad" : "_" + pn) +
101 (kHasLogitsSoftCap ? "_logits" : "_nlogits" ) + (BiasEnum == BlockAttentionBiasEnum::NO_BIAS ? _SS_("_nbias") : (_SS_("_") + BlockAttentionBiasEnumToStr<BiasEnum>::name)) +
102 (kHasMask ? "_" + _SS_(FmhaMask::name) : "_nmask") + (kStoreLSE ? "_lse" : "_nlse" ) + (kHasDropout ? "_dropout" : "_ndropout" ) + (kDoFp8StaticQuant ? "_squant" : "_nsquant" );
#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
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_dropout.hpp:26
Definition: block_position_encoding.hpp:137
Definition: fmha_batch_prefill_kernel.hpp:312
ck_tile::index_t kv_head_idx
Definition: fmha_batch_prefill_kernel.hpp:315
ck_tile::index_t qo_head_idx
Definition: fmha_batch_prefill_kernel.hpp:314
ck_tile::index_t batch_idx
Definition: fmha_batch_prefill_kernel.hpp:313
Definition: fmha_batch_prefill_kernel.hpp:192
ck_tile::index_t alibi_slope_stride
Definition: fmha_batch_prefill_kernel.hpp:195
const void * alibi_slope_ptr
Definition: fmha_batch_prefill_kernel.hpp:194
Definition: fmha_batch_prefill_kernel.hpp:187
ck_tile::index_t batch_stride_bias
Definition: fmha_batch_prefill_kernel.hpp:188
Definition: fmha_batch_prefill_kernel.hpp:268
ck_tile::index_t batch_stride_randval
Definition: fmha_batch_prefill_kernel.hpp:269
Definition: fmha_batch_prefill_kernel.hpp:284
ck_tile::index_t batch_stride_o
Definition: fmha_batch_prefill_kernel.hpp:288
ck_tile::index_t batch_stride_v
Definition: fmha_batch_prefill_kernel.hpp:287
ck_tile::index_t batch_stride_q
Definition: fmha_batch_prefill_kernel.hpp:285
ck_tile::index_t batch_stride_k
Definition: fmha_batch_prefill_kernel.hpp:286
Definition: fmha_batch_prefill_kernel.hpp:180
ck_tile::index_t nhead_stride_bias
Definition: fmha_batch_prefill_kernel.hpp:183
ck_tile::index_t stride_bias
Definition: fmha_batch_prefill_kernel.hpp:182
const void * bias_ptr
Definition: fmha_batch_prefill_kernel.hpp:181
Definition: fmha_batch_prefill_kernel.hpp:233
ck_tile::index_t stride_randval
Definition: fmha_batch_prefill_kernel.hpp:263
void init_dropout(float p_drop, const uint64_t *seed_ptr, const uint64_t *offset_ptr)
Definition: fmha_batch_prefill_kernel.hpp:246
ck_tile::index_t nhead_stride_randval
Definition: fmha_batch_prefill_kernel.hpp:264
void init_dropout(float p_drop, uint64_t seed, uint64_t offset)
Definition: fmha_batch_prefill_kernel.hpp:234
void * rand_val_ptr
Definition: fmha_batch_prefill_kernel.hpp:261
float rp_undrop
Definition: fmha_batch_prefill_kernel.hpp:258
bool is_store_randval
Definition: fmha_batch_prefill_kernel.hpp:260
uint8_t p_undrop_in_uint8_t
Definition: fmha_batch_prefill_kernel.hpp:259
Definition: fmha_batch_prefill_kernel.hpp:118
ck_tile::index_t stride_q
Definition: fmha_batch_prefill_kernel.hpp:146
ck_tile::index_t stride_v
Definition: fmha_batch_prefill_kernel.hpp:148
int32_t num_total_pages
Definition: fmha_batch_prefill_kernel.hpp:134
float scale_s
Definition: fmha_batch_prefill_kernel.hpp:144
ck_tile::index_t seqlen_q
Definition: fmha_batch_prefill_kernel.hpp:124
ck_tile::index_t stride_k
Definition: fmha_batch_prefill_kernel.hpp:147
ck_tile::index_t nhead_stride_o
Definition: fmha_batch_prefill_kernel.hpp:154
ck_tile::index_t nhead_stride_k
Definition: fmha_batch_prefill_kernel.hpp:152
ck_tile::index_t nhead_ratio_qk
Definition: fmha_batch_prefill_kernel.hpp:132
ck_tile::index_t nhead_stride_v
Definition: fmha_batch_prefill_kernel.hpp:153
ck_tile::index_t nhead_stride_q
Definition: fmha_batch_prefill_kernel.hpp:151
const int32_t * kv_page_indices
Definition: fmha_batch_prefill_kernel.hpp:136
const void * v_ptr
Definition: fmha_batch_prefill_kernel.hpp:121
const int32_t * kv_indptr
Definition: fmha_batch_prefill_kernel.hpp:135
void * o_ptr
Definition: fmha_batch_prefill_kernel.hpp:122
ck_tile::index_t seqlen_k
Definition: fmha_batch_prefill_kernel.hpp:125
ck_tile::index_t stride_o
Definition: fmha_batch_prefill_kernel.hpp:149
ck_tile::index_t hdim_v
Definition: fmha_batch_prefill_kernel.hpp:127
ck_tile::index_t num_head_q
Definition: fmha_batch_prefill_kernel.hpp:129
static constexpr ck_tile::index_t page_block_size
Definition: fmha_batch_prefill_kernel.hpp:141
const void * k_ptr
Definition: fmha_batch_prefill_kernel.hpp:120
ck_tile::index_t hdim_q
Definition: fmha_batch_prefill_kernel.hpp:126
const void * q_ptr
Definition: fmha_batch_prefill_kernel.hpp:119
Definition: fmha_batch_prefill_kernel.hpp:212
ck_tile::index_t batch_stride_lse
Definition: fmha_batch_prefill_kernel.hpp:215
ck_tile::index_t nhead_stride_lse
Definition: fmha_batch_prefill_kernel.hpp:214
void * lse_ptr
Definition: fmha_batch_prefill_kernel.hpp:213
Definition: fmha_batch_prefill_kernel.hpp:219
bool is_drop_seed_offset_from_host
Definition: fmha_batch_prefill_kernel.hpp:229
ValueOrPointer< uint64_t > drop_seed
Definition: fmha_batch_prefill_kernel.hpp:227
ValueOrPointer< uint64_t > drop_offset
Definition: fmha_batch_prefill_kernel.hpp:228
Definition: fmha_batch_prefill_kernel.hpp:111
Definition: fmha_batch_prefill_kernel.hpp:206
float scale_p
Definition: fmha_batch_prefill_kernel.hpp:207
float scale_o
Definition: fmha_batch_prefill_kernel.hpp:208
Definition: fmha_batch_prefill_kernel.hpp:303
ck_tile::index_t batch_stride_v
Definition: fmha_batch_prefill_kernel.hpp:306
ck_tile::index_t batch_stride_k
Definition: fmha_batch_prefill_kernel.hpp:305
const int32_t * seqstart_q_ptr
Definition: fmha_batch_prefill_kernel.hpp:304
Definition: fmha_batch_prefill_kernel.hpp:158
float logits_soft_cap_rcp
Definition: fmha_batch_prefill_kernel.hpp:176
void init_logits_soft_cap(float logits_soft_cap_)
Definition: fmha_batch_prefill_kernel.hpp:161
FmhaFwdLogitsSoftCapKargs()=default
float logits_soft_cap
Definition: fmha_batch_prefill_kernel.hpp:175
Definition: fmha_batch_prefill_kernel.hpp:199
ck_tile::index_t window_size_right
Definition: fmha_batch_prefill_kernel.hpp:201
ck_tile::index_t window_size_left
Definition: fmha_batch_prefill_kernel.hpp:201
ck_tile::GenericAttentionMaskEnum mask_type
Definition: fmha_batch_prefill_kernel.hpp:202
Definition: fmha_batch_prefill_kernel.hpp:64
Definition: fmha_batch_prefill_kernel.hpp:26
static constexpr CK_TILE_DEVICE auto GetTileIndex(const Kargs &kargs)
Definition: fmha_batch_prefill_kernel.hpp:636
static constexpr bool kIsGroupMode
Definition: fmha_batch_prefill_kernel.hpp:47
static constexpr ck_tile::index_t kBlockPerCu
Definition: fmha_batch_prefill_kernel.hpp:31
ck_tile::remove_cvref_t< typename FmhaPipeline::VDataType > VDataType
Definition: fmha_batch_prefill_kernel.hpp:37
ck_tile::remove_cvref_t< FmhaPipeline_ > FmhaPipeline
Definition: fmha_batch_prefill_kernel.hpp:27
ck_tile::remove_cvref_t< typename FmhaPipeline::KDataType > KDataType
Definition: fmha_batch_prefill_kernel.hpp:36
static constexpr bool kPadSeqLenQ
Definition: fmha_batch_prefill_kernel.hpp:48
ck_tile::remove_cvref_t< typename FmhaPipeline::RandValOutputDataType > RandValOutputDataType
Definition: fmha_batch_prefill_kernel.hpp:40
static constexpr bool kDoFp8StaticQuant
Definition: fmha_batch_prefill_kernel.hpp:56
static constexpr bool kPadHeadDimV
Definition: fmha_batch_prefill_kernel.hpp:51
ck_tile::remove_cvref_t< typename FmhaPipeline::LSEDataType > LSEDataType
Definition: fmha_batch_prefill_kernel.hpp:41
ck_tile::remove_cvref_t< typename FmhaPipeline::QDataType > QDataType
Definition: fmha_batch_prefill_kernel.hpp:35
static constexpr bool kHasMask
Definition: fmha_batch_prefill_kernel.hpp:59
static CK_TILE_HOST std::string GetName()
Definition: fmha_batch_prefill_kernel.hpp:72
ck_tile::remove_cvref_t< typename FmhaPipeline::BiasDataType > BiasDataType
Definition: fmha_batch_prefill_kernel.hpp:38
ck_tile::remove_cvref_t< typename FmhaPipeline::FmhaMask > FmhaMask
Definition: fmha_batch_prefill_kernel.hpp:58
static constexpr bool kPadSeqLenK
Definition: fmha_batch_prefill_kernel.hpp:49
static constexpr bool kHasLogitsSoftCap
Definition: fmha_batch_prefill_kernel.hpp:52
static constexpr bool kHasDropout
Definition: fmha_batch_prefill_kernel.hpp:55
static constexpr bool kStoreLSE
Definition: fmha_batch_prefill_kernel.hpp:54
ck_tile::remove_cvref_t< typename FmhaPipeline::SaccDataType > SaccDataType
Definition: fmha_batch_prefill_kernel.hpp:43
static constexpr CK_TILE_HOST_DEVICE ck_tile::index_t GetSmemSize()
Definition: fmha_batch_prefill_kernel.hpp:697
static constexpr ck_tile::index_t kBlockSize
Definition: fmha_batch_prefill_kernel.hpp:29
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 *rand_val_ptr, void *lse_ptr, void *o_ptr, const void *seqstart_q_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, int32_t num_total_pages, const void *kv_indptr, const void *kv_page_indices, 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_randval, 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_randval, 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, float p_drop, bool s_randval, std::variant< std::pair< uint64_t, uint64_t >, std::pair< const void *, const void * >> drop_seed_offset)
Definition: fmha_batch_prefill_kernel.hpp:472
static constexpr auto BiasEnum
Definition: fmha_batch_prefill_kernel.hpp:53
ck_tile::remove_cvref_t< typename FmhaPipeline::VLayout > VLayout
Definition: fmha_batch_prefill_kernel.hpp:45
ck_tile::remove_cvref_t< typename FmhaPipeline::ODataType > ODataType
Definition: fmha_batch_prefill_kernel.hpp:42
ck_tile::remove_cvref_t< typename FmhaPipeline::AttentionVariant > AttentionVariant
Definition: fmha_batch_prefill_kernel.hpp:57
static constexpr bool kUseAsyncCopy
Definition: fmha_batch_prefill_kernel.hpp:61
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 *rand_val_ptr, void *lse_ptr, void *o_ptr, ck_tile::index_t seqlen_q, 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, int32_t num_total_pages, const void *kv_indptr, const void *kv_page_indices, 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_randval, 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_randval, 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_randval, 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, float p_drop, bool s_randval, std::variant< std::pair< uint64_t, uint64_t >, std::pair< const void *, const void * >> drop_seed_offset)
Definition: fmha_batch_prefill_kernel.hpp:320
static constexpr ck_tile::index_t kBlockPerCuInput
Definition: fmha_batch_prefill_kernel.hpp:33
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_)
Definition: fmha_batch_prefill_kernel.hpp:613
static constexpr bool kPadHeadDimQ
Definition: fmha_batch_prefill_kernel.hpp:50
static constexpr CK_TILE_HOST auto BlockSize()
Definition: fmha_batch_prefill_kernel.hpp:695
CK_TILE_DEVICE void operator()(Kargs kargs) const
Definition: fmha_batch_prefill_kernel.hpp:702
std::conditional_t< kIsGroupMode, FmhaFwdGroupModeKargs, FmhaFwdBatchModeKargs > Kargs
Definition: fmha_batch_prefill_kernel.hpp:309
ck_tile::remove_cvref_t< EpiloguePipeline_ > EpiloguePipeline
Definition: fmha_batch_prefill_kernel.hpp:28
Definition: variants.hpp:63
Definition: block_dropout.hpp:12
Definition: variants.hpp:51
Definition: integral_constant.hpp:13
Definition: functional.hpp:86
Definition: coordinate_transform.hpp:1392
Definition: unary_element_function.hpp:56
Definition: math.hpp:28
Definition: sequence.hpp:49
Definition: fmha_batch_prefill_kernel.hpp:222
T val
Definition: fmha_batch_prefill_kernel.hpp:223
const T * ptr
Definition: fmha_batch_prefill_kernel.hpp:224