BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy > Struct Template Reference

BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP&lt; Problem, Policy &gt; Struct Template Reference#

Composable Kernel: ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy > Struct Template Reference

#include <block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr_iglp.hpp>

Public Types

using QDataType = remove_cvref_t<typename Problem::QDataType>
using KDataType = remove_cvref_t<typename Problem::KDataType>
using VDataType = remove_cvref_t<typename Problem::VDataType>
using GemmDataType = remove_cvref_t<typename Problem::GemmDataType>
using BiasDataType = remove_cvref_t<typename Problem::BiasDataType>
using LSEDataType = remove_cvref_t<typename Problem::LSEDataType>
using AccDataType = remove_cvref_t<typename Problem::AccDataType>
using DDataType = remove_cvref_t<typename Problem::DDataType>
using RandValOutputDataType = remove_cvref_t<typename Problem::RandValOutputDataType>
using ODataType = remove_cvref_t<typename Problem::ODataType>
using OGradDataType = remove_cvref_t<typename Problem::OGradDataType>
using QGradDataType = remove_cvref_t<typename Problem::QGradDataType>
using KGradDataType = remove_cvref_t<typename Problem::KGradDataType>
using VGradDataType = remove_cvref_t<typename Problem::VGradDataType>
using BiasGradDataType = remove_cvref_t<typename Problem::BiasGradDataType>
using FmhaMask = remove_cvref_t<typename Problem::FmhaMask>
using FmhaDropout = remove_cvref_t<typename Problem::FmhaDropout>
using HotLoopScheduler = typename Policy::template HotLoopScheduler<Problem>
using BlockFmhaShape = remove_cvref_t<typename Problem::BlockFmhaShape>

Public Member Functions

template<typename QDramBlockWindowTmp, typename KDramBlockWindowTmp, typename VDramBlockWindowTmp, typename BiasDramBlockWindowTmp, typename RandValDramBlockWindowTmp, typename OGradDramBlockWindowTmp, typename LSEDramBlockWindowTmp, typename DDramBlockWindowTmp, typename QGradDramBlockWindowTmp, typename BiasGradDramBlockWindowTmp, typename PositionEncoding>
CK_TILE_HOST_DEVICE auto operator() (void *smem_ptr, const QDramBlockWindowTmp &q_dram_block_window_tmp, const KDramBlockWindowTmp &k_dram_block_window_tmp, const VDramBlockWindowTmp &v_dram_block_window_tmp, const BiasDramBlockWindowTmp &bias_dram_block_window_tmp, const RandValDramBlockWindowTmp &randval_dram_block_window_tmp, const OGradDramBlockWindowTmp &do_dram_block_window_tmp, const LSEDramBlockWindowTmp &lse_dram_block_window_tmp, const DDramBlockWindowTmp &d_dram_block_window_tmp, const QGradDramBlockWindowTmp &dq_dram_block_window_tmp, const BiasGradDramBlockWindowTmp &dbias_dram_block_window_tmp, FmhaMask mask, PositionEncoding position_encoding, float raw_scale, float scale, float rp_undrop, float scale_rp_undrop, FmhaDropout &dropout) const

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSize ()

Static Public Attributes

static constexpr index_t kBlockPerCu = Problem::kBlockPerCu
static constexpr index_t kBlockSize = Problem::kBlockSize
static constexpr index_t kM0 = BlockFmhaShape::kM0
static constexpr index_t kN0 = BlockFmhaShape::kN0
static constexpr index_t kK0 = BlockFmhaShape::kK0
static constexpr index_t kK1 = BlockFmhaShape::kK1
static constexpr index_t kK2 = BlockFmhaShape::kK2
static constexpr index_t kK3 = BlockFmhaShape::kK3
static constexpr index_t kK4 = BlockFmhaShape::kK4
static constexpr index_t kQKHeaddim = BlockFmhaShape::kQKHeaddim
static constexpr index_t kVHeaddim = BlockFmhaShape::kVHeaddim
static constexpr bool kIsGroupMode = Problem::kIsGroupMode
static constexpr index_t kPadHeadDimQ = Problem::kPadHeadDimQ
static constexpr index_t kPadHeadDimV = Problem::kPadHeadDimV
static constexpr auto BiasEnum = Problem::BiasEnum
static constexpr bool kHasBiasGrad = Problem::kHasBiasGrad
static constexpr bool kIsDeterministic = Problem::kIsDeterministic
static constexpr bool kUseTrLoad = Problem::kUseTrLoad
static constexpr index_t kAlignmentQ
static constexpr index_t kAlignmentK
static constexpr index_t kAlignmentV
static constexpr index_t kAlignmentOGrad
static constexpr index_t kAlignmentQGrad = 1
static constexpr index_t kAlignmentKGrad
static constexpr index_t kAlignmentVGrad
static constexpr index_t kAlignmentBias = 1
static constexpr const char * name = "kr_ktr_vr_iglp"

Member Typedef Documentation

◆ AccDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::AccDataType = remove_cvref_t<typename Problem::AccDataType>

◆ BiasDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::BiasDataType = remove_cvref_t<typename Problem::BiasDataType>

◆ BiasGradDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::BiasGradDataType = remove_cvref_t<typename Problem::BiasGradDataType>

◆ BlockFmhaShape

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::BlockFmhaShape = remove_cvref_t<typename Problem::BlockFmhaShape>

◆ DDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::DDataType = remove_cvref_t<typename Problem::DDataType>

◆ FmhaDropout

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::FmhaDropout = remove_cvref_t<typename Problem::FmhaDropout>

◆ FmhaMask

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::FmhaMask = remove_cvref_t<typename Problem::FmhaMask>

◆ GemmDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::GemmDataType = remove_cvref_t<typename Problem::GemmDataType>

◆ HotLoopScheduler

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::HotLoopScheduler = typename Policy::template HotLoopScheduler<Problem>

◆ KDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::KDataType = remove_cvref_t<typename Problem::KDataType>

◆ KGradDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::KGradDataType = remove_cvref_t<typename Problem::KGradDataType>

◆ LSEDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::LSEDataType = remove_cvref_t<typename Problem::LSEDataType>

◆ ODataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::ODataType = remove_cvref_t<typename Problem::ODataType>

◆ OGradDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::OGradDataType = remove_cvref_t<typename Problem::OGradDataType>

◆ QDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::QDataType = remove_cvref_t<typename Problem::QDataType>

◆ QGradDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::QGradDataType = remove_cvref_t<typename Problem::QGradDataType>

◆ RandValOutputDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::RandValOutputDataType = remove_cvref_t<typename Problem::RandValOutputDataType>

◆ VDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::VDataType = remove_cvref_t<typename Problem::VDataType>

◆ VGradDataType

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
using ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::VGradDataType = remove_cvref_t<typename Problem::VGradDataType>

Member Function Documentation

◆ GetSmemSize()

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::GetSmemSize ( )
inlinestaticconstexpr

◆ operator()()

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
template<typename QDramBlockWindowTmp, typename KDramBlockWindowTmp, typename VDramBlockWindowTmp, typename BiasDramBlockWindowTmp, typename RandValDramBlockWindowTmp, typename OGradDramBlockWindowTmp, typename LSEDramBlockWindowTmp, typename DDramBlockWindowTmp, typename QGradDramBlockWindowTmp, typename BiasGradDramBlockWindowTmp, typename PositionEncoding>
CK_TILE_HOST_DEVICE auto ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::operator() ( void * smem_ptr,
const QDramBlockWindowTmp & q_dram_block_window_tmp,
const KDramBlockWindowTmp & k_dram_block_window_tmp,
const VDramBlockWindowTmp & v_dram_block_window_tmp,
const BiasDramBlockWindowTmp & bias_dram_block_window_tmp,
const RandValDramBlockWindowTmp & randval_dram_block_window_tmp,
const OGradDramBlockWindowTmp & do_dram_block_window_tmp,
const LSEDramBlockWindowTmp & lse_dram_block_window_tmp,
const DDramBlockWindowTmp & d_dram_block_window_tmp,
const QGradDramBlockWindowTmp & dq_dram_block_window_tmp,
const BiasGradDramBlockWindowTmp & dbias_dram_block_window_tmp,
FmhaMask mask,
PositionEncoding position_encoding,
float raw_scale,
float scale,
float rp_undrop,
float scale_rp_undrop,
FmhaDropout & dropout ) const
inline

Member Data Documentation

◆ BiasEnum

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
auto ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::BiasEnum = Problem::BiasEnum
staticconstexpr

◆ kAlignmentBias

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kAlignmentBias = 1
staticconstexpr

◆ kAlignmentK

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kAlignmentK
staticconstexpr
Initial value:
=
kPadHeadDimQ ? kPadHeadDimQ : Policy::template GetAlignmentK<Problem>()
ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVR::kPadHeadDimQ
static constexpr index_t kPadHeadDimQ
Definition block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr.hpp:52

◆ kAlignmentKGrad

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kAlignmentKGrad
staticconstexpr
Initial value:
=
kPadHeadDimQ ? kPadHeadDimQ : Policy::template GetAlignmentKGrad<Problem>()

◆ kAlignmentOGrad

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kAlignmentOGrad
staticconstexpr
Initial value:
=
kPadHeadDimV ? kPadHeadDimV : Policy::template GetAlignmentOGrad<Problem>()
ck_tile::BlockFmhaBwdOGradDotO::kPadHeadDimV
static constexpr bool kPadHeadDimV
Definition block_fmha_bwd_dot_do_o.hpp:24

◆ kAlignmentQ

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kAlignmentQ
staticconstexpr
Initial value:
=
kPadHeadDimQ ? kPadHeadDimQ : Policy::template GetAlignmentQ<Problem>()

◆ kAlignmentQGrad

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kAlignmentQGrad = 1
staticconstexpr

◆ kAlignmentV

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kAlignmentV
staticconstexpr
Initial value:
=
kPadHeadDimV ? kPadHeadDimV : Policy::template GetAlignmentV<Problem>()
ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVR::kPadHeadDimV
static constexpr index_t kPadHeadDimV
Definition block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr.hpp:53

◆ kAlignmentVGrad

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kAlignmentVGrad
staticconstexpr
Initial value:
=
kPadHeadDimV ? kPadHeadDimV : Policy::template GetAlignmentVGrad<Problem>()

◆ kBlockPerCu

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kBlockPerCu = Problem::kBlockPerCu
staticconstexpr

◆ kBlockSize

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kBlockSize = Problem::kBlockSize
staticconstexpr

◆ kHasBiasGrad

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
bool ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kHasBiasGrad = Problem::kHasBiasGrad
staticconstexpr

◆ kIsDeterministic

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
bool ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kIsDeterministic = Problem::kIsDeterministic
staticconstexpr

◆ kIsGroupMode

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
bool ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kIsGroupMode = Problem::kIsGroupMode
staticconstexpr

◆ kK0

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kK0 = BlockFmhaShape::kK0
staticconstexpr

◆ kK1

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kK1 = BlockFmhaShape::kK1
staticconstexpr

◆ kK2

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kK2 = BlockFmhaShape::kK2
staticconstexpr

◆ kK3

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kK3 = BlockFmhaShape::kK3
staticconstexpr

◆ kK4

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kK4 = BlockFmhaShape::kK4
staticconstexpr

◆ kM0

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kM0 = BlockFmhaShape::kM0
staticconstexpr

◆ kN0

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kN0 = BlockFmhaShape::kN0
staticconstexpr

◆ kPadHeadDimQ

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kPadHeadDimQ = Problem::kPadHeadDimQ
staticconstexpr

◆ kPadHeadDimV

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kPadHeadDimV = Problem::kPadHeadDimV
staticconstexpr

◆ kQKHeaddim

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kQKHeaddim = BlockFmhaShape::kQKHeaddim
staticconstexpr

◆ kUseTrLoad

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
bool ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kUseTrLoad = Problem::kUseTrLoad
staticconstexpr

◆ kVHeaddim

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::kVHeaddim = BlockFmhaShape::kVHeaddim
staticconstexpr

◆ name

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
const char* ck_tile::BlockFmhaBwdDQDKDVPipelineKRKTRVRIGLP< Problem, Policy >::name = "kr_ktr_vr_iglp"
staticconstexpr
The documentation for this struct was generated from the following file: