BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ > Struct Template Reference

BlockFmhaPipelineProblem&lt; QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ > Struct Template Reference
Public Types | Static Public Attributes | List of all members
ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, 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 BiasDataType = remove_cvref_t<BiasDataType_>
using RandValOutputDataType = remove_cvref_t<RandValOutputDataType_>
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 AttentionVariant = remove_cvref_t<AttentionVariant_>
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 kUseTrLoad = kUseTrLoad_
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 kHasLogitsSoftCap = Traits::kHasLogitsSoftCap
static constexpr bool kSkipMinSeqlenQ = Traits::kSkipMinSeqlenQ
static constexpr auto BiasEnum = Traits::BiasEnum
static constexpr bool kStoreLSE = Traits::kStoreLSE
static constexpr bool kHasDropout = Traits::kHasDropout
static constexpr bool kDoFp8StaticQuant = Traits::kDoFp8StaticQuant
static constexpr index_t kBlockPerCu = Traits::kBlockPerCu

Member Typedef Documentation

◆ AttentionVariant

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::AttentionVariant = remove_cvref_t<AttentionVariant_>

◆ BiasDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::BiasDataType = remove_cvref_t<BiasDataType_>

◆ BlockFmhaShape

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

◆ FmhaMask

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

◆ KDataType

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

◆ LSEDataType

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

◆ OaccDataType

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

◆ ODataType

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

◆ PDataType

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

◆ QDataType

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

◆ RandValOutputDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::RandValOutputDataType = remove_cvref_t<RandValOutputDataType_>

◆ SaccDataType

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

◆ SMPLComputeDataType

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

◆ Traits

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

◆ VDataType

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

Member Data Documentation

◆ BiasEnum

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
auto ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::BiasEnum = Traits::BiasEnum
staticconstexpr

◆ kBlockPerCu

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

◆ kBlockSize

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

◆ kDoFp8StaticQuant

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::kDoFp8StaticQuant = Traits::kDoFp8StaticQuant
staticconstexpr

◆ kHasDropout

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::kHasDropout = Traits::kHasDropout
staticconstexpr

◆ kHasLogitsSoftCap

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::kHasLogitsSoftCap = Traits::kHasLogitsSoftCap
staticconstexpr

◆ kIsGroupMode

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

◆ kNumGemm0Warps

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

◆ kNumGemm1Warps

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

◆ kPadHeadDimQ

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

◆ kPadHeadDimV

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

◆ kPadSeqLenK

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

◆ kPadSeqLenQ

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

◆ kSkipMinSeqlenQ

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::kSkipMinSeqlenQ = Traits::kSkipMinSeqlenQ
staticconstexpr

◆ kStoreLSE

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::kStoreLSE = Traits::kStoreLSE
staticconstexpr

◆ kUseTrLoad

template<typename QDataType_, typename KDataType_, typename VDataType_, typename SaccDataType_, typename SMPLComputeDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename LSEDataType_, typename PDataType_, typename OaccDataType_, typename ODataType_, typename BlockFmhaShape_, bool kIsGroupMode_, typename AttentionVariant_, typename FmhaMask_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, AttentionVariant_, FmhaMask_, kUseTrLoad_, Traits_ >::kUseTrLoad = kUseTrLoad_
staticconstexpr
The documentation for this struct was generated from the following file: