BlockFmhaBwdPipelineTrLoadDefaultPolicy Struct Reference

• ck_tile
• BlockFmhaBwdPipelineTrLoadDefaultPolicy

Classes | Static Public Member Functions | Static Public Attributes | List of all members

ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy Struct Reference

#include <block_fmha_bwd_pipeline_trload_default_policy.hpp>

Classes
class	HotLoopScheduler

Static Public Member Functions
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	GetQKBlockGemm ()
template<typename Problem>
static CK_TILE_HOST_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, typename T>
static CK_TILE_HOST_DEVICE constexpr auto	GetAlignmentX () noexcept
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 T>
static CK_TILE_HOST_DEVICE constexpr auto	GetTransposedAlignmentX () noexcept
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	GetTransposedAlignmentQ () noexcept
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 T, typename TensorView>
static CK_TILE_HOST_DEVICE constexpr auto	TransformXDramTensorView (const TensorView &naive_view)
template<typename T, typename... TD_TS>
static CK_TILE_HOST_DEVICE constexpr auto	TransformXDramDescriptor (const tensor_descriptor< TD_TS... > &from_desc)
template<typename Problem, typename T, index_t RowsPerBlock, index_t ColsPerBlock>
static CK_TILE_HOST_DEVICE constexpr auto	MakeXDramTileDistribution ()
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>
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	MakeKRegBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeVRegBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeKTRegBlockDescriptor ()
template<typename T, index_t MNPerBlock, index_t KPerBlock>
static CK_TILE_HOST_DEVICE constexpr auto	MakeXLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeKLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeVLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeQLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeOGradLdsWriteBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeBiasLdsBlockDescriptor ()
template<typename Problem, bool Transposed = false>
static CK_TILE_HOST_DEVICE constexpr auto	MakeSGradLdsBlockDescriptor ()
template<typename T, index_t MNPerBlock, index_t KPerBlock>
static CK_TILE_HOST_DEVICE constexpr auto	MakeXLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeKLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeVLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeQLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeOGradLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeQRegSliceBlockDescriptor ()
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>
static CK_TILE_HOST_DEVICE constexpr auto	MakeLSEDLdsReadBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeOGradRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_DEVICE constexpr auto	MakeOGradTRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_DEVICE constexpr auto	MakePTRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeSGradRegSliceBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeShuffledBiasTileDistribution ()
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	GetSmemSizeK ()
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	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
static constexpr index_t	WarpAlignmentBytes = 128

Member Function Documentation

◆ GetAlignmentBias()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentBias ( )

inlinestaticconstexpr

◆ GetAlignmentK()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentK ( )

inlinestaticconstexpr

◆ GetAlignmentKGrad()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentKGrad ( )

inlinestaticconstexpr

◆ GetAlignmentO()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentO ( )

inlinestaticconstexpr

◆ GetAlignmentOGrad()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentOGrad ( )

inlinestaticconstexpr

◆ GetAlignmentPostQGrad()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentPostQGrad ( )

inlinestaticconstexpr

◆ GetAlignmentPostQGradAcc()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentPostQGradAcc ( )

inlinestaticconstexpr

◆ GetAlignmentQ()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentQ ( )

inlinestaticconstexpr

◆ GetAlignmentV()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentV ( )

inlinestaticconstexpr

◆ GetAlignmentVGrad()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentVGrad ( )

inlinestaticconstexpr

◆ GetAlignmentX()

template<typename Problem, typename T>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetAlignmentX ( )

inlinestaticconstexprnoexcept

◆ GetOGradVBlockGemm()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetOGradVBlockGemm ( )

inlinestaticconstexpr

◆ GetPTOGradTBlockGemm()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetPTOGradTBlockGemm ( )

inlinestaticconstexpr

◆ GetQKBlockGemm()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetQKBlockGemm ( )

inlinestaticconstexpr

◆ GetSGradKTBlockGemm()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSGradKTBlockGemm ( )

inlinestaticconstexpr

◆ GetSGradTQTBlockGemm()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSGradTQTBlockGemm ( )

inlinestaticconstexpr

◆ GetSmemSize()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSize ( )

inlinestaticconstexpr

◆ GetSmemSizeBias()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSizeBias ( )

inlinestaticconstexpr

◆ GetSmemSizeD()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSizeD ( )

inlinestaticconstexpr

◆ GetSmemSizeK()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSizeK ( )

inlinestaticconstexpr

◆ GetSmemSizeLSE()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSizeLSE ( )

inlinestaticconstexpr

◆ GetSmemSizeOGrad()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSizeOGrad ( )

inlinestaticconstexpr

◆ GetSmemSizeQ()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSizeQ ( )

inlinestaticconstexpr

◆ GetSmemSizeSGrad()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSizeSGrad ( )

inlinestaticconstexpr

◆ GetSmemSizeV()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetSmemSizeV ( )

inlinestaticconstexpr

◆ GetTransposedAlignmentBias()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetTransposedAlignmentBias ( )

inlinestaticconstexpr

◆ GetTransposedAlignmentOGrad()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetTransposedAlignmentOGrad ( )

inlinestaticconstexpr

◆ GetTransposedAlignmentQ()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetTransposedAlignmentQ ( )

inlinestaticconstexprnoexcept

◆ GetTransposedAlignmentX()

template<typename T>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::GetTransposedAlignmentX ( )

inlinestaticconstexprnoexcept

◆ MakeBiasLdsBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeBiasLdsBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeBiasSTileDistribution()

template<typename BlockGemm>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeBiasSTileDistribution ( )

inlinestaticconstexpr

◆ MakeBiasTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeBiasTileDistribution ( )

inlinestaticconstexpr

◆ MakeKDramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeKDramTileDistribution ( )

inlinestaticconstexpr

◆ MakeKLdsReadBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeKLdsReadBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeKLdsWriteBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeKLdsWriteBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeKRegBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeKRegBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeKTRegBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeKTRegBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeLSEDDramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeLSEDDramTileDistribution ( )

inlinestaticconstexpr

◆ MakeLSEDLdsReadBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeLSEDLdsReadBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeLSEDLdsWriteBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeLSEDLdsWriteBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeOGradDramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeOGradDramTileDistribution ( )

inlinestaticconstexpr

◆ MakeOGradLdsReadBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeOGradLdsReadBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeOGradLdsWriteBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeOGradLdsWriteBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeOGradRegSliceBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeOGradRegSliceBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeOGradTRegSliceBlockDescriptor()

template<typename Problem>

CK_TILE_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeOGradTRegSliceBlockDescriptor ( )

inlinestaticconstexpr

◆ MakePostQGradAccDramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakePostQGradAccDramTileDistribution ( )

inlinestaticconstexpr

◆ MakePostQGradDramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakePostQGradDramTileDistribution ( )

inlinestaticconstexpr

◆ MakePreODramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakePreODramTileDistribution ( )

inlinestaticconstexpr

◆ MakePreOGradDramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakePreOGradDramTileDistribution ( )

inlinestaticconstexpr

◆ MakePreXDramTileDistribution()

template<typename DataType, index_t MPerBlock, index_t KPerBlock>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakePreXDramTileDistribution ( )

inlinestaticconstexpr

◆ MakePTRegSliceBlockDescriptor()

template<typename Problem>

CK_TILE_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakePTRegSliceBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeQDramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeQDramTileDistribution ( )

inlinestaticconstexpr

◆ MakeQLdsReadBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeQLdsReadBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeQLdsWriteBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeQLdsWriteBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeQRegSliceBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeQRegSliceBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeQTRegSliceBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeQTRegSliceBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeSGradLdsBlockDescriptor()

template<typename Problem, bool Transposed = false>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeSGradLdsBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeSGradRegSliceBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeSGradRegSliceBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeSGradTRegSliceBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeSGradTRegSliceBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeShuffledBiasTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeShuffledBiasTileDistribution ( )

inlinestaticconstexpr

◆ MakeVDramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeVDramTileDistribution ( )

inlinestaticconstexpr

◆ MakeVLdsReadBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeVLdsReadBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeVLdsWriteBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeVLdsWriteBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeVRegBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeVRegBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeXDramTileDistribution()

template<typename Problem, typename T, index_t RowsPerBlock, index_t ColsPerBlock>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeXDramTileDistribution ( )

inlinestaticconstexpr

◆ MakeXLdsReadBlockDescriptor()

template<typename T, index_t MNPerBlock, index_t KPerBlock>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeXLdsReadBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeXLdsWriteBlockDescriptor()

template<typename T, index_t MNPerBlock, index_t KPerBlock>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::MakeXLdsWriteBlockDescriptor ( )

inlinestaticconstexpr

◆ TransformXDramDescriptor()

template<typename T, typename... TD_TS>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::TransformXDramDescriptor ( const tensor_descriptor< TD_TS... > & from_desc )

inlinestaticconstexpr

◆ TransformXDramTensorView()

template<typename T, typename TensorView>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::TransformXDramTensorView ( const TensorView & naive_view )

inlinestaticconstexpr

Member Data Documentation

◆ WarpAlignmentBytes

index_t ck_tile::BlockFmhaBwdPipelineTrLoadDefaultPolicy::WarpAlignmentBytes = 128

staticconstexpr

---

The documentation for this struct was generated from the following file:

• block_fmha_bwd_pipeline_trload_default_policy.hpp