FmhaFwdAppendKVKernel< FmhaPipeline_ > Struct Template Reference

FmhaFwdAppendKVKernel&lt; FmhaPipeline_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ > Struct Template Reference
Classes | Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ > Struct Template Reference

#include <fmha_fwd_appendkv_kernel.hpp>

Classes

struct  t2s
struct  t2s< float >
struct  t2s< ck_tile::fp16_t >
struct  t2s< ck_tile::bf16_t >
struct  t2s< ck_tile::fp8_t >
struct  t2s< ck_tile::bf8_t >
struct  EmptyKargs
struct  BasicKargs
struct  RoPEKargs
struct  PageBlockTableKargs
struct  CacheBatchIdxKargs
struct  Kargs

Public Types

using FmhaPipeline = ck_tile::remove_cvref_t<FmhaPipeline_>
using QDataType = ck_tile::remove_cvref_t<typename FmhaPipeline::QDataType>
using KDataType = ck_tile::remove_cvref_t<typename FmhaPipeline::KDataType>
using VDataType = ck_tile::remove_cvref_t<typename FmhaPipeline::VDataType>
using VLayout = ck_tile::remove_cvref_t<typename FmhaPipeline::VLayout>

Public Member Functions

CK_TILE_DEVICE void operator() (Kargs kargs) const

Static Public Member Functions

static CK_TILE_HOST std::string GetName ()
static CK_TILE_HOST constexpr Kargs MakeKargs (void *q_ptr, void *k_ptr, const void *knew_ptr, void *v_ptr, const void *vnew_ptr, ck_tile::index_t seqlen_q, const void *seqlen_k_ptr, ck_tile::index_t seqlen_knew, 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 *rotary_cos_ptr, const void *rotary_sin_ptr, ck_tile::index_t rotary_dim, bool has_mask, 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, ck_tile::index_t stride_q, ck_tile::index_t stride_k, ck_tile::index_t stride_knew, ck_tile::index_t stride_v, ck_tile::index_t stride_vnew, ck_tile::index_t nhead_stride_q, ck_tile::index_t nhead_stride_k, ck_tile::index_t nhead_stride_knew, ck_tile::index_t nhead_stride_v, ck_tile::index_t nhead_stride_vnew, ck_tile::index_t batch_stride_q, ck_tile::index_t batch_stride_k, ck_tile::index_t batch_stride_knew, ck_tile::index_t batch_stride_v, ck_tile::index_t batch_stride_vnew)
static CK_TILE_HOST constexpr auto GridSize (ck_tile::index_t batch_size, ck_tile::index_t nhead, ck_tile::index_t seqlen_q, ck_tile::index_t seqlen_knew)
static CK_TILE_DEVICE constexpr auto GetTileIndex (const Kargs &)
static CK_TILE_HOST dim3 BlockSize ()

Static Public Attributes

static constexpr ck_tile::index_t kBlockSize = FmhaPipeline::kBlockSize
static constexpr ck_tile::index_t kBlockPerCu = FmhaPipeline::kBlockPerCu
static constexpr ck_tile::index_t kBlockPerCuInput = FmhaPipeline::Problem::kBlockPerCu
static constexpr bool kApplyRoPE = FmhaPipeline::RotaryEnum != RotaryEmbeddingEnum::NONE
static constexpr bool kIsPagedKV = FmhaPipeline::kIsPagedKV
static constexpr bool kPadSeqLenQ = FmhaPipeline::kPadSeqLenQ
static constexpr bool kPadSeqLenK = FmhaPipeline::kPadSeqLenK
static constexpr bool kPadHeadDimQ = FmhaPipeline::kPadHeadDimQ
static constexpr bool kPadHeadDimV = FmhaPipeline::kPadHeadDimV

Member Typedef Documentation

◆ FmhaPipeline

template<typename FmhaPipeline_>
using ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::FmhaPipeline = ck_tile::remove_cvref_t<FmhaPipeline_>

◆ KDataType

template<typename FmhaPipeline_>
using ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::KDataType = ck_tile::remove_cvref_t<typename FmhaPipeline::KDataType>

◆ QDataType

template<typename FmhaPipeline_>
using ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::QDataType = ck_tile::remove_cvref_t<typename FmhaPipeline::QDataType>

◆ VDataType

template<typename FmhaPipeline_>
using ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::VDataType = ck_tile::remove_cvref_t<typename FmhaPipeline::VDataType>

◆ VLayout

template<typename FmhaPipeline_>
using ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::VLayout = ck_tile::remove_cvref_t<typename FmhaPipeline::VLayout>

Member Function Documentation

◆ BlockSize()

template<typename FmhaPipeline_>
CK_TILE_HOST dim3 ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::BlockSize ( )
inlinestatic

◆ GetName()

template<typename FmhaPipeline_>
CK_TILE_HOST std::string ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::GetName ( )
inlinestatic

◆ GetTileIndex()

template<typename FmhaPipeline_>
CK_TILE_DEVICE constexpr auto ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::GetTileIndex ( const Kargs & )
inlinestaticconstexpr

◆ GridSize()

template<typename FmhaPipeline_>
CK_TILE_HOST constexpr auto ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::GridSize ( ck_tile::index_t batch_size,
ck_tile::index_t nhead,
ck_tile::index_t seqlen_q,
ck_tile::index_t seqlen_knew )
inlinestaticconstexpr

◆ MakeKargs()

template<typename FmhaPipeline_>
CK_TILE_HOST constexpr Kargs ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::MakeKargs ( void * q_ptr,
void * k_ptr,
const void * knew_ptr,
void * v_ptr,
const void * vnew_ptr,
ck_tile::index_t seqlen_q,
const void * seqlen_k_ptr,
ck_tile::index_t seqlen_knew,
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 * rotary_cos_ptr,
const void * rotary_sin_ptr,
ck_tile::index_t rotary_dim,
bool has_mask,
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,
ck_tile::index_t stride_q,
ck_tile::index_t stride_k,
ck_tile::index_t stride_knew,
ck_tile::index_t stride_v,
ck_tile::index_t stride_vnew,
ck_tile::index_t nhead_stride_q,
ck_tile::index_t nhead_stride_k,
ck_tile::index_t nhead_stride_knew,
ck_tile::index_t nhead_stride_v,
ck_tile::index_t nhead_stride_vnew,
ck_tile::index_t batch_stride_q,
ck_tile::index_t batch_stride_k,
ck_tile::index_t batch_stride_knew,
ck_tile::index_t batch_stride_v,
ck_tile::index_t batch_stride_vnew )
inlinestaticconstexpr

◆ operator()()

template<typename FmhaPipeline_>
CK_TILE_DEVICE void ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::operator() ( Kargs kargs) const
inline

Member Data Documentation

◆ kApplyRoPE

template<typename FmhaPipeline_>
bool ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kApplyRoPE = FmhaPipeline::RotaryEnum != RotaryEmbeddingEnum::NONE
staticconstexpr

◆ kBlockPerCu

template<typename FmhaPipeline_>
ck_tile::index_t ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kBlockPerCu = FmhaPipeline::kBlockPerCu
staticconstexpr

◆ kBlockPerCuInput

template<typename FmhaPipeline_>
ck_tile::index_t ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kBlockPerCuInput = FmhaPipeline::Problem::kBlockPerCu
staticconstexpr

◆ kBlockSize

template<typename FmhaPipeline_>
ck_tile::index_t ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kBlockSize = FmhaPipeline::kBlockSize
staticconstexpr

◆ kIsPagedKV

template<typename FmhaPipeline_>
bool ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kIsPagedKV = FmhaPipeline::kIsPagedKV
staticconstexpr

◆ kPadHeadDimQ

template<typename FmhaPipeline_>
bool ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kPadHeadDimQ = FmhaPipeline::kPadHeadDimQ
staticconstexpr

◆ kPadHeadDimV

template<typename FmhaPipeline_>
bool ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kPadHeadDimV = FmhaPipeline::kPadHeadDimV
staticconstexpr

◆ kPadSeqLenK

template<typename FmhaPipeline_>
bool ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kPadSeqLenK = FmhaPipeline::kPadSeqLenK
staticconstexpr

◆ kPadSeqLenQ

template<typename FmhaPipeline_>
bool ck_tile::FmhaFwdAppendKVKernel< FmhaPipeline_ >::kPadSeqLenQ = FmhaPipeline::kPadSeqLenQ
staticconstexpr
The documentation for this struct was generated from the following file: