BlockFmhaBwdConvertQGrad< Problem, Policy > Struct Template Reference

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

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

#include <block_fmha_bwd_convert_dq.hpp>

Public Types

using AccDataType = remove_cvref_t<typename Problem::AccDataType>
using QGradDataType = remove_cvref_t<typename Problem::QGradDataType>

Public Member Functions

template<typename QGradAccDramBlockWindowTmp, typename QGradDramBlockWindowTmp>
CK_TILE_HOST_DEVICE void operator() (const QGradAccDramBlockWindowTmp &dq_acc_dram_block_window_tmp, QGradDramBlockWindowTmp &dq_dram_block_window_tmp) const
template<typename QGradAccDramBlockWindowTmp, typename QGradDramBlockWindowTmp>
CK_TILE_HOST_DEVICE void operator() (const QGradAccDramBlockWindowTmp &dq_acc_dram_block_window_tmp, QGradDramBlockWindowTmp &dq_dram_block_window_tmp, index_t nsplits) const

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSize ()

Static Public Attributes

static constexpr index_t kM0 = Problem::kM0
static constexpr index_t kN0 = Problem::kN0
static constexpr index_t kBlockPerCu = Problem::kBlockPerCu
static constexpr index_t kBlockSize = Problem::kBlockSize
static constexpr index_t kQKHeaddim = Problem::kQKHeaddim
static constexpr bool kIsGroupMode = Problem::kIsGroupMode
static constexpr bool kPadSeqLenQ = Problem::kPadSeqLenQ
static constexpr bool kPadHeadDimQ = Problem::kPadHeadDimQ
static constexpr bool kIsDeterministic = Problem::kIsDeterministic
static constexpr index_t kAlignmentQGradAcc
static constexpr index_t kAlignmentQGrad

Member Typedef Documentation

◆ AccDataType

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

◆ QGradDataType

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

Member Function Documentation

◆ GetSmemSize()

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

◆ operator()() [1/2]

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
template<typename QGradAccDramBlockWindowTmp, typename QGradDramBlockWindowTmp>
CK_TILE_HOST_DEVICE void ck_tile::BlockFmhaBwdConvertQGrad< Problem, Policy >::operator() ( const QGradAccDramBlockWindowTmp & dq_acc_dram_block_window_tmp,
QGradDramBlockWindowTmp & dq_dram_block_window_tmp ) const
inline

◆ operator()() [2/2]

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
template<typename QGradAccDramBlockWindowTmp, typename QGradDramBlockWindowTmp>
CK_TILE_HOST_DEVICE void ck_tile::BlockFmhaBwdConvertQGrad< Problem, Policy >::operator() ( const QGradAccDramBlockWindowTmp & dq_acc_dram_block_window_tmp,
QGradDramBlockWindowTmp & dq_dram_block_window_tmp,
index_t nsplits ) const
inline

Member Data Documentation

◆ kAlignmentQGrad

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdConvertQGrad< Problem, Policy >::kAlignmentQGrad
staticconstexpr
Initial value:
=
kPadHeadDimQ ? 1 : Policy::template GetAlignmentPostQGrad<Problem>()
ck_tile::BlockFmhaBwdConvertQGrad::kPadHeadDimQ
static constexpr bool kPadHeadDimQ
Definition block_fmha_bwd_convert_dq.hpp:26

◆ kAlignmentQGradAcc

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdConvertQGrad< Problem, Policy >::kAlignmentQGradAcc
staticconstexpr
Initial value:
=
kPadHeadDimQ ? 1 : Policy::template GetAlignmentPostQGradAcc<Problem>()

◆ kBlockPerCu

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

◆ kBlockSize

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

◆ kIsDeterministic

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

◆ kIsGroupMode

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

◆ kM0

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

◆ kN0

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

◆ kPadHeadDimQ

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
bool ck_tile::BlockFmhaBwdConvertQGrad< Problem, Policy >::kPadHeadDimQ = Problem::kPadHeadDimQ
staticconstexpr

◆ kPadSeqLenQ

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
bool ck_tile::BlockFmhaBwdConvertQGrad< Problem, Policy >::kPadSeqLenQ = Problem::kPadSeqLenQ
staticconstexpr

◆ kQKHeaddim

template<typename Problem, typename Policy = BlockFmhaBwdPipelineDefaultPolicy>
index_t ck_tile::BlockFmhaBwdConvertQGrad< Problem, Policy >::kQKHeaddim = Problem::kQKHeaddim
staticconstexpr
The documentation for this struct was generated from the following file: