BlockFmhaBwdPipelineDefaultPolicy Struct Reference

BlockFmhaBwdPipelineDefaultPolicy Struct Reference#

Composable Kernel: ck_tile::BlockFmhaBwdPipelineDefaultPolicy Struct Reference
Classes | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::BlockFmhaBwdPipelineDefaultPolicy Struct Reference

#include <block_fmha_bwd_pipeline_default_policy.hpp>

Classes

struct  HotLoopScheduler

Static Public Member Functions

template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetQKBlockGemm ()
template<typename Problem>
static CK_TILE_DEVICE constexpr auto GetPTOGradTBlockGemm ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetOGradVBlockGemm ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSGradTQTBlockGemm ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSGradKTBlockGemm ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentQ ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentK ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentV ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentO ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentOGrad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentBias ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentKGrad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentVGrad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetTransposedAlignmentQ ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetTransposedAlignmentK ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetTransposedAlignmentOGrad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetTransposedAlignmentBias ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentPostQGradAcc ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentPostQGrad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeKDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeVDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeQDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeOGradDramTileDistribution ()
template<typename Problem, typename BlockGemm>
static CK_TILE_HOST_DEVICE constexpr auto MakeLSEDDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeBiasTileDistribution ()
template<typename DataType, index_t MPerBlock, index_t KPerBlock>
static CK_TILE_HOST_DEVICE constexpr auto MakePreXDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakePreODramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakePreOGradDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakePostQGradAccDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakePostQGradDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackQ ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackQT ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackK ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackKT ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackV ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackBias ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackBiasT ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackOGrad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackOGradT ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackSGrad ()
template<index_t KIter, index_t MNPerBlock, index_t KPerSubBlock, index_t KPack>
static CK_TILE_HOST_DEVICE constexpr auto MakeXLdsBlockDescriptor ()
template<index_t MNPerBlock, index_t KPerBlock, index_t KPack>
static CK_TILE_HOST_DEVICE constexpr auto MakeXLdsBlockDescriptor ()
template<typename Problem, index_t MNPerBlock, index_t KPerBlock, index_t KPack, index_t KPackT>
static CK_TILE_HOST_DEVICE constexpr auto MakeXTLdsBlockDescriptor ()
template<typename Problem, index_t MNIter, index_t MNPerSubBlock, index_t KPerBlock, index_t KPack, index_t KPackT>
static CK_TILE_HOST_DEVICE constexpr auto MakeXTLdsBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeKLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeKRegBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeVLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeVRegBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledKRegWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledKLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeKTLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeKTRegBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeQLdsBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeQRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledQRegWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledQLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeQTLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeQTRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeSGradTRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeLSEDLdsWriteBlockDescriptor ()
template<typename Problem, typename BlockGemm>
static CK_TILE_HOST_DEVICE constexpr auto MakeLSEDLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeOGradLdsBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeOGradRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledOGradRegWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledOGradLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeOGradTLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeOGradTRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakePTRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeSGradLdsBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeSGradRegSliceBlockDescriptor ()
template<typename Problem, typename PTOutTensor, typename PInTensor>
static CK_TILE_DEVICE constexpr void PTFromGemm0CToGemm1A (PTOutTensor &pt_out, const PInTensor &p_in)
template<typename Problem, typename SGradTOutTensor, typename SGradInTensor>
static CK_TILE_DEVICE constexpr void SGradTFromGemm2CToGemm3A (SGradTOutTensor &dst_out, const SGradInTensor &ds_in)
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledBiasTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeBiasLdsBlockDescriptor ()
template<typename BlockGemm>
static CK_TILE_HOST_DEVICE constexpr auto MakeBiasSTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeQ ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeQT ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeK ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeKT ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeLSE ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeD ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeV ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeOGrad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeOGradT ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeSGrad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeBias ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()

Static Public Attributes

template<index_t ndim>
static constexpr auto swap_last2

Member Function Documentation

◆ GetAlignmentBias()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentBias ( )
inlinestaticconstexpr

◆ GetAlignmentK()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentK ( )
inlinestaticconstexpr

◆ GetAlignmentKGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentKGrad ( )
inlinestaticconstexpr

◆ GetAlignmentO()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentO ( )
inlinestaticconstexpr

◆ GetAlignmentOGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentOGrad ( )
inlinestaticconstexpr

◆ GetAlignmentPostQGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentPostQGrad ( )
inlinestaticconstexpr

◆ GetAlignmentPostQGradAcc()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentPostQGradAcc ( )
inlinestaticconstexpr

◆ GetAlignmentQ()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentQ ( )
inlinestaticconstexpr

◆ GetAlignmentV()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentV ( )
inlinestaticconstexpr

◆ GetAlignmentVGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetAlignmentVGrad ( )
inlinestaticconstexpr

◆ GetOGradVBlockGemm()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetOGradVBlockGemm ( )
inlinestaticconstexpr

◆ GetPTOGradTBlockGemm()

template<typename Problem>
CK_TILE_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetPTOGradTBlockGemm ( )
inlinestaticconstexpr

◆ GetQKBlockGemm()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetQKBlockGemm ( )
inlinestaticconstexpr

◆ GetSGradKTBlockGemm()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSGradKTBlockGemm ( )
inlinestaticconstexpr

◆ GetSGradTQTBlockGemm()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSGradTQTBlockGemm ( )
inlinestaticconstexpr

◆ GetSmemKPackBias()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackBias ( )
inlinestaticconstexpr

◆ GetSmemKPackBiasT()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackBiasT ( )
inlinestaticconstexpr

◆ GetSmemKPackK()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackK ( )
inlinestaticconstexpr

◆ GetSmemKPackKT()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackKT ( )
inlinestaticconstexpr

◆ GetSmemKPackOGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackOGrad ( )
inlinestaticconstexpr

◆ GetSmemKPackOGradT()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackOGradT ( )
inlinestaticconstexpr

◆ GetSmemKPackQ()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackQ ( )
inlinestaticconstexpr

◆ GetSmemKPackQT()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackQT ( )
inlinestaticconstexpr

◆ GetSmemKPackSGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackSGrad ( )
inlinestaticconstexpr

◆ GetSmemKPackV()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemKPackV ( )
inlinestaticconstexpr

◆ GetSmemSize()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSize ( )
inlinestaticconstexpr

◆ GetSmemSizeBias()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeBias ( )
inlinestaticconstexpr

◆ GetSmemSizeD()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeD ( )
inlinestaticconstexpr

◆ GetSmemSizeK()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeK ( )
inlinestaticconstexpr

◆ GetSmemSizeKT()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeKT ( )
inlinestaticconstexpr

◆ GetSmemSizeLSE()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeLSE ( )
inlinestaticconstexpr

◆ GetSmemSizeOGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeOGrad ( )
inlinestaticconstexpr

◆ GetSmemSizeOGradT()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeOGradT ( )
inlinestaticconstexpr

◆ GetSmemSizeQ()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeQ ( )
inlinestaticconstexpr

◆ GetSmemSizeQT()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeQT ( )
inlinestaticconstexpr

◆ GetSmemSizeSGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeSGrad ( )
inlinestaticconstexpr

◆ GetSmemSizeV()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetSmemSizeV ( )
inlinestaticconstexpr

◆ GetTransposedAlignmentBias()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetTransposedAlignmentBias ( )
inlinestaticconstexpr

◆ GetTransposedAlignmentK()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetTransposedAlignmentK ( )
inlinestaticconstexpr

◆ GetTransposedAlignmentOGrad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetTransposedAlignmentOGrad ( )
inlinestaticconstexpr

◆ GetTransposedAlignmentQ()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::GetTransposedAlignmentQ ( )
inlinestaticconstexpr

◆ MakeBiasLdsBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeBiasLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeBiasSTileDistribution()

template<typename BlockGemm>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeBiasSTileDistribution ( )
inlinestaticconstexpr

◆ MakeBiasTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeBiasTileDistribution ( )
inlinestaticconstexpr

◆ MakeKDramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeKDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeKLdsWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeKLdsWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeKRegBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeKRegBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeKTLdsReadBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeKTLdsReadBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeKTRegBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeKTRegBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeLSEDDramTileDistribution()

template<typename Problem, typename BlockGemm>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeLSEDDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeLSEDLdsReadBlockDescriptor()

template<typename Problem, typename BlockGemm>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeLSEDLdsReadBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeLSEDLdsWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeLSEDLdsWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeOGradDramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeOGradDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeOGradLdsBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeOGradLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeOGradRegSliceBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeOGradRegSliceBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeOGradTLdsReadBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeOGradTLdsReadBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeOGradTRegSliceBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeOGradTRegSliceBlockDescriptor ( )
inlinestaticconstexpr

◆ MakePostQGradAccDramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakePostQGradAccDramTileDistribution ( )
inlinestaticconstexpr

◆ MakePostQGradDramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakePostQGradDramTileDistribution ( )
inlinestaticconstexpr

◆ MakePreODramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakePreODramTileDistribution ( )
inlinestaticconstexpr

◆ MakePreOGradDramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakePreOGradDramTileDistribution ( )
inlinestaticconstexpr

◆ MakePreXDramTileDistribution()

template<typename DataType, index_t MPerBlock, index_t KPerBlock>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakePreXDramTileDistribution ( )
inlinestaticconstexpr

◆ MakePTRegSliceBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakePTRegSliceBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeQDramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeQDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeQLdsBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeQLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeQRegSliceBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeQRegSliceBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeQTLdsReadBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeQTLdsReadBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeQTRegSliceBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeQTRegSliceBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeSGradLdsBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeSGradLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeSGradRegSliceBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeSGradRegSliceBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeSGradTRegSliceBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeSGradTRegSliceBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeShuffledBiasTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeShuffledBiasTileDistribution ( )
inlinestaticconstexpr

◆ MakeShuffledKLdsWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeShuffledKLdsWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeShuffledKRegWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeShuffledKRegWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeShuffledOGradLdsWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeShuffledOGradLdsWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeShuffledOGradRegWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeShuffledOGradRegWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeShuffledQLdsWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeShuffledQLdsWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeShuffledQRegWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeShuffledQRegWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeVDramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeVDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeVLdsWriteBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeVLdsWriteBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeVRegBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeVRegBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeXLdsBlockDescriptor() [1/2]

template<index_t MNPerBlock, index_t KPerBlock, index_t KPack>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeXLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeXLdsBlockDescriptor() [2/2]

template<index_t KIter, index_t MNPerBlock, index_t KPerSubBlock, index_t KPack>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeXLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeXTLdsBlockDescriptor() [1/2]

template<typename Problem, index_t MNIter, index_t MNPerSubBlock, index_t KPerBlock, index_t KPack, index_t KPackT>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeXTLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeXTLdsBlockDescriptor() [2/2]

template<typename Problem, index_t MNPerBlock, index_t KPerBlock, index_t KPack, index_t KPackT>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::MakeXTLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ PTFromGemm0CToGemm1A()

template<typename Problem, typename PTOutTensor, typename PInTensor>
CK_TILE_DEVICE constexpr void ck_tile::BlockFmhaBwdPipelineDefaultPolicy::PTFromGemm0CToGemm1A ( PTOutTensor & pt_out,
const PInTensor & p_in )
inlinestaticconstexpr

◆ SGradTFromGemm2CToGemm3A()

template<typename Problem, typename SGradTOutTensor, typename SGradInTensor>
CK_TILE_DEVICE constexpr void ck_tile::BlockFmhaBwdPipelineDefaultPolicy::SGradTFromGemm2CToGemm3A ( SGradTOutTensor & dst_out,
const SGradInTensor & ds_in )
inlinestaticconstexpr

Member Data Documentation

◆ swap_last2

template<index_t ndim>
auto ck_tile::BlockFmhaBwdPipelineDefaultPolicy::swap_last2
staticconstexpr
Initial value:
= generate_sequence_v2(
[](auto i) {
return number < i == ndim - 2 ? ndim - 1 : i == ndim - 1 ? ndim - 2 : i > {};
},
number<ndim>{})
ck_tile::generate_sequence_v2
CK_TILE_HOST_DEVICE constexpr auto generate_sequence_v2(F &&f, number< N >)
Definition tile/core/container/sequence.hpp:1045
ck_tile::number
constant< v > number
Definition tile/core/numeric/integral_constant.hpp:37
The documentation for this struct was generated from the following file: