BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ > Struct Template Reference

BlockFmhaFwdAppendKVPipelineProblem&lt; QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ > Struct Template Reference
Public Types | Static Public Attributes | List of all members
ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ > Struct Template Reference

#include <block_fmha_pipeline_problem.hpp>

Public Types

using QDataType = remove_cvref_t<QDataType_>
using KDataType = remove_cvref_t<KDataType_>
using VDataType = remove_cvref_t<VDataType_>
using Traits = remove_cvref_t<Traits_>
using VLayout

Static Public Attributes

static constexpr index_t kBlockSize = 256
static constexpr index_t kM0 = kM0_
static constexpr index_t kN0 = kN0_
static constexpr index_t kK0 = kK0_
static constexpr index_t kN1 = kN1_
static constexpr auto RotaryEnum = RotaryEnum_
static constexpr bool kIsPagedKV = kIsPagedKV_
static constexpr bool kPadSeqLenQ = Traits::kPadSeqLenQ
static constexpr bool kPadSeqLenK = Traits::kPadSeqLenK
static constexpr bool kPadHeadDimQ = Traits::kPadHeadDimQ
static constexpr bool kPadHeadDimV = Traits::kPadHeadDimV
static constexpr index_t kBlockPerCu = Traits::kBlockPerCu

Member Typedef Documentation

◆ KDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
using ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::KDataType = remove_cvref_t<KDataType_>

◆ QDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
using ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::QDataType = remove_cvref_t<QDataType_>

◆ Traits

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
using ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::Traits = remove_cvref_t<Traits_>

◆ VDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
using ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::VDataType = remove_cvref_t<VDataType_>

◆ VLayout

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
using ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::VLayout
Initial value:
std::conditional_t<kIsVLayoutRowMajor_,
ck_tile::tensor_layout::gemm::RowMajor,
ck_tile::tensor_layout::gemm::ColumnMajor>
ck_tile::tensor_layout::gemm::ColumnMajor
Definition tile/ops/common/tensor_layout.hpp:22
ck_tile::tensor_layout::gemm::RowMajor
Definition tile/ops/common/tensor_layout.hpp:17

Member Data Documentation

◆ kBlockPerCu

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
index_t ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kBlockPerCu = Traits::kBlockPerCu
staticconstexpr

◆ kBlockSize

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
index_t ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kBlockSize = 256
staticconstexpr

◆ kIsPagedKV

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
bool ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kIsPagedKV = kIsPagedKV_
staticconstexpr

◆ kK0

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
index_t ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kK0 = kK0_
staticconstexpr

◆ kM0

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
index_t ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kM0 = kM0_
staticconstexpr

◆ kN0

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
index_t ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kN0 = kN0_
staticconstexpr

◆ kN1

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
index_t ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kN1 = kN1_
staticconstexpr

◆ kPadHeadDimQ

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
bool ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kPadHeadDimQ = Traits::kPadHeadDimQ
staticconstexpr

◆ kPadHeadDimV

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
bool ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kPadHeadDimV = Traits::kPadHeadDimV
staticconstexpr

◆ kPadSeqLenK

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
bool ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kPadSeqLenK = Traits::kPadSeqLenK
staticconstexpr

◆ kPadSeqLenQ

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
bool ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::kPadSeqLenQ = Traits::kPadSeqLenQ
staticconstexpr

◆ RotaryEnum

template<typename QDataType_, typename KDataType_, typename VDataType_, index_t kM0_, index_t kN0_, index_t kK0_, index_t kN1_, bool kIsVLayoutRowMajor_, RotaryEmbeddingEnum RotaryEnum_, bool kIsPagedKV_, typename Traits_>
auto ck_tile::BlockFmhaFwdAppendKVPipelineProblem< QDataType_, KDataType_, VDataType_, kM0_, kN0_, kK0_, kN1_, kIsVLayoutRowMajor_, RotaryEnum_, kIsPagedKV_, Traits_ >::RotaryEnum = RotaryEnum_
staticconstexpr
The documentation for this struct was generated from the following file: