BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ > Struct Template Reference

BlockFmhaFwdV3PipelineProblem&lt; QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ > Struct Template Reference
Public Types | Static Public Attributes | List of all members
ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 SaccDataType = remove_cvref_t<SaccDataType_>
using SMPLComputeDataType = remove_cvref_t<SMPLComputeDataType_>
using LSEDataType = remove_cvref_t<LSEDataType_>
using PDataType = remove_cvref_t<PDataType_>
using OaccDataType = remove_cvref_t<OaccDataType_>
using ODataType = remove_cvref_t<ODataType_>
using BlockFmhaShape = remove_cvref_t<BlockFmhaShape_>
using FmhaMask = remove_cvref_t<FmhaMask_>
using Traits = remove_cvref_t<Traits_>

Static Public Attributes

static constexpr index_t kNumGemm0Warps = BlockFmhaShape::NumGemm0Warps
static constexpr index_t kNumGemm1Warps = BlockFmhaShape::NumGemm1Warps
static constexpr index_t kBlockSize = BlockFmhaShape::NumWarps * get_warp_size()
static constexpr bool kIsGroupMode = kIsGroupMode_
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 bool kStoreLSE = Traits::kStoreLSE
static constexpr index_t kBlockPerCu = Traits::kBlockPerCu

Member Typedef Documentation

◆ BlockFmhaShape

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::BlockFmhaShape = remove_cvref_t<BlockFmhaShape_>

◆ FmhaMask

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::FmhaMask = remove_cvref_t<FmhaMask_>

◆ KDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::KDataType = remove_cvref_t<KDataType_>

◆ LSEDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::LSEDataType = remove_cvref_t<LSEDataType_>

◆ OaccDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::OaccDataType = remove_cvref_t<OaccDataType_>

◆ ODataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::ODataType = remove_cvref_t<ODataType_>

◆ PDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::PDataType = remove_cvref_t<PDataType_>

◆ QDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::QDataType = remove_cvref_t<QDataType_>

◆ SaccDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::SaccDataType = remove_cvref_t<SaccDataType_>

◆ SMPLComputeDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::SMPLComputeDataType = remove_cvref_t<SMPLComputeDataType_>

◆ Traits

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::Traits = remove_cvref_t<Traits_>

◆ VDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
using ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::VDataType = remove_cvref_t<VDataType_>

Member Data Documentation

◆ kBlockPerCu

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
index_t ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kBlockPerCu = Traits::kBlockPerCu
staticconstexpr

◆ kBlockSize

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
index_t ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kBlockSize = BlockFmhaShape::NumWarps * get_warp_size()
staticconstexpr

◆ kIsGroupMode

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
bool ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kIsGroupMode = kIsGroupMode_
staticconstexpr

◆ kNumGemm0Warps

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
index_t ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kNumGemm0Warps = BlockFmhaShape::NumGemm0Warps
staticconstexpr

◆ kNumGemm1Warps

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
index_t ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kNumGemm1Warps = BlockFmhaShape::NumGemm1Warps
staticconstexpr

◆ kPadHeadDimQ

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
bool ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kPadHeadDimQ = Traits::kPadHeadDimQ
staticconstexpr

◆ kPadHeadDimV

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
bool ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kPadHeadDimV = Traits::kPadHeadDimV
staticconstexpr

◆ kPadSeqLenK

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
bool ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kPadSeqLenK = Traits::kPadSeqLenK
staticconstexpr

◆ kPadSeqLenQ

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
bool ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kPadSeqLenQ = Traits::kPadSeqLenQ
staticconstexpr

◆ kStoreLSE

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename FmhaMask_, typename Traits_>
bool ck_tile::BlockFmhaFwdV3PipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kStoreLSE = Traits::kStoreLSE
staticconstexpr
The documentation for this struct was generated from the following file: