BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ > Struct Template Reference

BlockFmhaBwdPipelineProblem&lt; QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ > Struct Template Reference
Public Types | Static Public Attributes | List of all members
ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ > Struct Template Reference

#include <block_fmha_bwd_pipeline_problem.hpp>

Public Types

using QDataType = remove_cvref_t<QDataType_>
using KDataType = remove_cvref_t<KDataType_>
using VDataType = remove_cvref_t<VDataType_>
using GemmDataType = remove_cvref_t<GemmDataType_>
using LSEDataType = remove_cvref_t<LSEDataType_>
using AccDataType = remove_cvref_t<AccDataType_>
using DDataType = remove_cvref_t<DDataType_>
using BiasDataType = remove_cvref_t<BiasDataType_>
using RandValOutputDataType = remove_cvref_t<RandValOutputDataType_>
using ODataType = remove_cvref_t<ODataType_>
using OGradDataType = remove_cvref_t<OGradDataType_>
using QGradDataType = remove_cvref_t<QGradDataType_>
using KGradDataType = remove_cvref_t<KGradDataType_>
using VGradDataType = remove_cvref_t<VGradDataType_>
using BiasGradDataType = remove_cvref_t<BiasGradDataType_>
using BlockFmhaShape = remove_cvref_t<BlockFmhaShape_>
using FmhaMask = remove_cvref_t<FmhaMask_>
using FmhaDropout = remove_cvref_t<FmhaDropout_>
using Traits = remove_cvref_t<Traits_>

Static Public Attributes

static constexpr index_t kBlockSize = BlockFmhaShape::NumWarps * get_warp_size()
static constexpr bool kIsGroupMode = kIsGroupMode_
static constexpr bool kIsDeterministic = kIsDeterministic_
static constexpr bool kUseTrLoad = kUseTrLoad_
static constexpr index_t kPadHeadDimQ = Traits::kPadHeadDimQ
static constexpr index_t kPadHeadDimV = Traits::kPadHeadDimV
static constexpr auto BiasEnum = Traits::BiasEnum
static constexpr bool kHasBiasGrad = Traits::kHasBiasGrad
static constexpr index_t kBlockPerCu = Traits::kBlockPerCu

Member Typedef Documentation

◆ AccDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::AccDataType = remove_cvref_t<AccDataType_>

◆ BiasDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::BiasDataType = remove_cvref_t<BiasDataType_>

◆ BiasGradDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::BiasGradDataType = remove_cvref_t<BiasGradDataType_>

◆ BlockFmhaShape

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::BlockFmhaShape = remove_cvref_t<BlockFmhaShape_>

◆ DDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::DDataType = remove_cvref_t<DDataType_>

◆ FmhaDropout

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::FmhaDropout = remove_cvref_t<FmhaDropout_>

◆ FmhaMask

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::FmhaMask = remove_cvref_t<FmhaMask_>

◆ GemmDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::GemmDataType = remove_cvref_t<GemmDataType_>

◆ KDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::KDataType = remove_cvref_t<KDataType_>

◆ KGradDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::KGradDataType = remove_cvref_t<KGradDataType_>

◆ LSEDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::LSEDataType = remove_cvref_t<LSEDataType_>

◆ ODataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::ODataType = remove_cvref_t<ODataType_>

◆ OGradDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::OGradDataType = remove_cvref_t<OGradDataType_>

◆ QDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::QDataType = remove_cvref_t<QDataType_>

◆ QGradDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::QGradDataType = remove_cvref_t<QGradDataType_>

◆ RandValOutputDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::RandValOutputDataType = remove_cvref_t<RandValOutputDataType_>

◆ Traits

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::Traits = remove_cvref_t<Traits_>

◆ VDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::VDataType = remove_cvref_t<VDataType_>

◆ VGradDataType

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
using ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::VGradDataType = remove_cvref_t<VGradDataType_>

Member Data Documentation

◆ BiasEnum

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
auto ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::BiasEnum = Traits::BiasEnum
staticconstexpr

◆ kBlockPerCu

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
index_t ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::kBlockPerCu = Traits::kBlockPerCu
staticconstexpr

◆ kBlockSize

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
index_t ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::kBlockSize = BlockFmhaShape::NumWarps * get_warp_size()
staticconstexpr

◆ kHasBiasGrad

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::kHasBiasGrad = Traits::kHasBiasGrad
staticconstexpr

◆ kIsDeterministic

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::kIsDeterministic = kIsDeterministic_
staticconstexpr

◆ kIsGroupMode

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::kIsGroupMode = kIsGroupMode_
staticconstexpr

◆ kPadHeadDimQ

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
index_t ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::kPadHeadDimQ = Traits::kPadHeadDimQ
staticconstexpr

◆ kPadHeadDimV

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
index_t ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::kPadHeadDimV = Traits::kPadHeadDimV
staticconstexpr

◆ kUseTrLoad

template<typename QDataType_, typename KDataType_, typename VDataType_, typename GemmDataType_, typename LSEDataType_, typename AccDataType_, typename DDataType_, typename BiasDataType_, typename RandValOutputDataType_, typename ODataType_, typename OGradDataType_, typename QGradDataType_, typename KGradDataType_, typename VGradDataType_, typename BiasGradDataType_, typename BlockFmhaShape_, bool kIsGroupMode_, bool kIsDeterministic_, typename FmhaMask_, typename FmhaDropout_, bool kUseTrLoad_, typename Traits_>
bool ck_tile::BlockFmhaBwdPipelineProblem< QDataType_, KDataType_, VDataType_, GemmDataType_, LSEDataType_, AccDataType_, DDataType_, BiasDataType_, RandValOutputDataType_, ODataType_, OGradDataType_, QGradDataType_, KGradDataType_, VGradDataType_, BiasGradDataType_, BlockFmhaShape_, kIsGroupMode_, kIsDeterministic_, FmhaMask_, FmhaDropout_, kUseTrLoad_, Traits_ >::kUseTrLoad = kUseTrLoad_
staticconstexpr
The documentation for this struct was generated from the following file: