BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy Struct Reference

BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy Struct Reference#

Composable Kernel: ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy Struct Reference
Public Types | Static Public Member Functions | List of all members
ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy Struct Reference

#include <block_fmha_pipeline_qr_ks_vs_async_trload_policy.hpp>

Inheritance diagram for ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy:
ck_tile::BlockFmhaPipelineQXKSVSCustomPolicy< true, false, 1, 1 > ck_tile::BlockFmhaPipelineQXCustomPolicy< QLoadOnce_ >

Public Types

using BasePolicy
Public Types inherited from ck_tile::BlockFmhaPipelineQXKSVSCustomPolicy< true, false, 1, 1 >
using QXPolicy

Static Public Member Functions

template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentQ ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentOacc ()
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, bool BypassLDS = false>
static CK_TILE_HOST_DEVICE constexpr auto MakeQDramTileDistribution ()
template<typename Problem, bool LoadOnce = false>
static CK_TILE_HOST_DEVICE constexpr auto MakeKDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeQRegTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackQ ()
template<typename Problem, bool Xor = false>
static CK_TILE_HOST_DEVICE constexpr auto MakeQLdsBlockDescriptor ()
template<typename Problem, bool LoadOnce = false, bool Xor = false>
static CK_TILE_HOST_DEVICE constexpr auto MakeKLdsBlockDescriptor ()
template<typename Problem, bool Xor = false>
static CK_TILE_HOST_DEVICE constexpr auto MakeVLdsBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetQKBlockGemm ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetPVBlockGemm ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeKRegTileDistribution ()
template<typename Problem>
static CK_TILE_DEVICE constexpr auto MakeVDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakePRegTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeVRegTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemNPackS ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeSLdsBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeSRegTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSizeQ ()
template<typename Problem, bool LoadOnce = false>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSizeK ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSizeV ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSizeS ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSize ()
Static Public Member Functions inherited from ck_tile::BlockFmhaPipelineQXKSVSCustomPolicy< true, false, 1, 1 >
static CK_TILE_HOST_DEVICE constexpr auto GetLdsBufferSequence ()
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackK ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentK ()
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackV ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentV ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentBias ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentO ()
static CK_TILE_HOST_DEVICE constexpr auto GetSingleSmemElementSpaceSize ()
static CK_TILE_HOST_DEVICE constexpr auto MakeKLdsBlockDescriptor ()
static CK_TILE_HOST_DEVICE constexpr auto MakeKLdsStoreBlockDescriptor (number< IBuf >=number< 0 >{})
static CK_TILE_HOST_DEVICE constexpr auto MakeKLdsLoadBlockDescriptor ()
static CK_TILE_HOST_DEVICE constexpr auto MakeVLdsBlockDescriptor ()
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSizeKV ()
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSize ()
static CK_TILE_HOST_DEVICE constexpr std::enable_if_t< std::is_convertible_v< decltype(Problem::kHasDropout), bool >, ck_tile::index_t > GetSmemSizeDropout (int)
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSizeDropout (...)
static CK_TILE_HOST_DEVICE constexpr auto MakeKDramTileDistribution ()
static CK_TILE_DEVICE constexpr auto MakeVDramTileDistribution ()
static CK_TILE_HOST_DEVICE constexpr auto MakeBiasDramTileDistribution ()
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledVRegBlockDescriptor ()
static CK_TILE_HOST_DEVICE constexpr auto GetKVBlockGemm ()

Additional Inherited Members

Static Public Attributes inherited from ck_tile::BlockFmhaPipelineQXKSVSCustomPolicy< true, false, 1, 1 >
static constexpr bool AsyncCopy
static constexpr index_t NumPrefetchK
static constexpr index_t NumPrefetchV
static constexpr index_t NumKVLdsBuffers

Member Typedef Documentation

◆ BasePolicy

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::BasePolicy
Initial value:
BlockFmhaPipelineQXKSVSCustomPolicy< true,
false,
1,
1>
ck_tile::BlockFmhaPipelineQXKSVSCustomPolicy
Definition block_fmha_pipeline_qx_ks_vs_custom_policy.hpp:266

Member Function Documentation

◆ GetAlignmentK()

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

◆ GetAlignmentOacc()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::GetAlignmentOacc ( )
inlinestaticconstexpr

◆ GetAlignmentQ()

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

◆ GetAlignmentV()

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

◆ GetPVBlockGemm()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::GetPVBlockGemm ( )
inlinestaticconstexpr

◆ GetQKBlockGemm()

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

◆ GetSmemKPackQ()

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

◆ GetSmemNPackS()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::GetSmemNPackS ( )
inlinestaticconstexpr

◆ GetSmemSize()

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

◆ GetSmemSizeK()

template<typename Problem, bool LoadOnce = false>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::GetSmemSizeK ( )
inlinestaticconstexpr

◆ GetSmemSizeQ()

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

◆ GetSmemSizeS()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::GetSmemSizeS ( )
inlinestaticconstexpr

◆ GetSmemSizeV()

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

◆ MakeKDramTileDistribution()

template<typename Problem, bool LoadOnce = false>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeKDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeKLdsBlockDescriptor()

template<typename Problem, bool LoadOnce = false, bool Xor = false>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeKLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeKRegTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeKRegTileDistribution ( )
inlinestaticconstexpr

◆ MakePRegTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakePRegTileDistribution ( )
inlinestaticconstexpr

◆ MakeQDramTileDistribution()

template<typename Problem, bool BypassLDS = false>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeQDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeQLdsBlockDescriptor()

template<typename Problem, bool Xor = false>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeQLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeQRegTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeQRegTileDistribution ( )
inlinestaticconstexpr

◆ MakeSLdsBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeSLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeSRegTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeSRegTileDistribution ( )
inlinestaticconstexpr

◆ MakeVDramTileDistribution()

template<typename Problem>
CK_TILE_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeVDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeVLdsBlockDescriptor()

template<typename Problem, bool Xor = false>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeVLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeVRegTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy::MakeVRegTileDistribution ( )
inlinestaticconstexpr
The documentation for this struct was generated from the following file: