BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ > Struct Template Reference

BlockFmhaSplitKVCombinePipelineProblem&lt; LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ > Struct Template Reference
Public Types | Static Public Attributes | List of all members
ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ > Struct Template Reference

#include <block_fmha_pipeline_problem.hpp>

Inheritance diagram for ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >:
ck_tile::BlockFmhaSplitKVCombinePipelineTileSizes< OaccDataType_, kN1_ >

Public Types

using BaseType = BlockFmhaSplitKVCombinePipelineTileSizes<OaccDataType_, kN1_>
using LSEDataType = remove_cvref_t<LSEDataType_>
using OaccDataType = remove_cvref_t<OaccDataType_>
using ODataType = remove_cvref_t<ODataType_>
using Traits = remove_cvref_t<Traits_>

Static Public Attributes

static constexpr index_t kHeadDimV = HeadDimV_
static constexpr bool kIsGroupMode = kIsGroupMode_
static constexpr bool kPadSeqLenQ = Traits::kPadSeqLenQ
static constexpr bool kPadHeadDimV = Traits::kPadHeadDimV
static constexpr bool kStoreLSE = Traits::kStoreLSE
static constexpr bool kDoFp8StaticQuant = Traits::kDoFp8StaticQuant
static constexpr index_t kBlockPerCu = Traits::kBlockPerCu
static constexpr index_t kMaxSplits = Traits::kMaxSplits
static constexpr index_t kNumWarps = 4
static constexpr index_t kBlockSize = kNumWarps * get_warp_size()
static constexpr index_t kM0
static constexpr index_t kN1
static constexpr index_t NThreads
Static Public Attributes inherited from ck_tile::BlockFmhaSplitKVCombinePipelineTileSizes< OaccDataType_, kN1_ >
static constexpr index_t MaxVectorSize = 16 / sizeof(OaccDataType_)
static constexpr index_t kN1 = kN1_
static constexpr index_t NThreads = kN1 / MaxVectorSize
static constexpr index_t kM0 = get_warp_size() / NThreads

Member Typedef Documentation

◆ BaseType

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
using ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::BaseType = BlockFmhaSplitKVCombinePipelineTileSizes<OaccDataType_, kN1_>

◆ LSEDataType

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
using ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::LSEDataType = remove_cvref_t<LSEDataType_>

◆ OaccDataType

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
using ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::OaccDataType = remove_cvref_t<OaccDataType_>

◆ ODataType

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
using ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::ODataType = remove_cvref_t<ODataType_>

◆ Traits

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
using ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::Traits = remove_cvref_t<Traits_>

Member Data Documentation

◆ kBlockPerCu

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
index_t ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kBlockPerCu = Traits::kBlockPerCu
staticconstexpr

◆ kBlockSize

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
index_t ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kBlockSize = kNumWarps * get_warp_size()
staticconstexpr

◆ kDoFp8StaticQuant

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
bool ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kDoFp8StaticQuant = Traits::kDoFp8StaticQuant
staticconstexpr

◆ kHeadDimV

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
index_t ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kHeadDimV = HeadDimV_
staticconstexpr

◆ kIsGroupMode

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
bool ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kIsGroupMode = kIsGroupMode_
staticconstexpr

◆ kM0

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
index_t ck_tile::BlockFmhaSplitKVCombinePipelineTileSizes< OaccDataType_, kN1_ >::kM0
staticconstexpr

◆ kMaxSplits

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
index_t ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kMaxSplits = Traits::kMaxSplits
staticconstexpr

◆ kN1

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
index_t ck_tile::BlockFmhaSplitKVCombinePipelineTileSizes< OaccDataType_, kN1_ >::kN1
staticconstexpr

◆ kNumWarps

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
index_t ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kNumWarps = 4
staticconstexpr

◆ kPadHeadDimV

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
bool ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kPadHeadDimV = Traits::kPadHeadDimV
staticconstexpr

◆ kPadSeqLenQ

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
bool ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kPadSeqLenQ = Traits::kPadSeqLenQ
staticconstexpr

◆ kStoreLSE

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
bool ck_tile::BlockFmhaSplitKVCombinePipelineProblem< LSEDataType_, OaccDataType_, ODataType_, HeadDimV_, kIsGroupMode_, kN1_, Traits_ >::kStoreLSE = Traits::kStoreLSE
staticconstexpr

◆ NThreads

template<typename LSEDataType_, typename OaccDataType_, typename ODataType_, index_t HeadDimV_, bool kIsGroupMode_, ck_tile::index_t kN1_, typename Traits_>
index_t ck_tile::BlockFmhaSplitKVCombinePipelineTileSizes< OaccDataType_, kN1_ >::NThreads
staticconstexpr
The documentation for this struct was generated from the following file: