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

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

Composable Kernel: ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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_, 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 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 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 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

◆ 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 FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
using ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 FmhaMask_ , typename Traits_ >
constexpr auto ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 FmhaMask_ , typename Traits_ >
constexpr index_t ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
constexpr index_t ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 FmhaMask_ , typename Traits_ >
constexpr bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 FmhaMask_ , typename Traits_ >
constexpr bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kHasDropout = Traits::kHasDropout
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 FmhaMask_ , typename Traits_ >
constexpr bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, 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 FmhaMask_ , typename Traits_ >
constexpr index_t ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
constexpr index_t ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
constexpr bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
constexpr bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
constexpr bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
constexpr bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, 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 BiasDataType_ , typename RandValOutputDataType_ , typename LSEDataType_ , typename PDataType_ , typename OaccDataType_ , typename ODataType_ , typename BlockFmhaShape_ , bool kIsGroupMode_, typename FmhaMask_ , typename Traits_ >
constexpr bool ck_tile::BlockFmhaPipelineProblem< QDataType_, KDataType_, VDataType_, SaccDataType_, SMPLComputeDataType_, BiasDataType_, RandValOutputDataType_, LSEDataType_, PDataType_, OaccDataType_, ODataType_, BlockFmhaShape_, kIsGroupMode_, FmhaMask_, Traits_ >::kStoreLSE = Traits::kStoreLSE
staticconstexpr
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_problem.hpp