BlockFmhaFwdSplitKVPipelineNWarpSShuffleQRKSVSDefaultPolicy Struct Reference

BlockFmhaFwdSplitKVPipelineNWarpSShuffleQRKSVSDefaultPolicy Struct Reference#

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

#include <block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs_default_policy.hpp>

Inheritance diagram for ck_tile::BlockFmhaFwdSplitKVPipelineNWarpSShuffleQRKSVSDefaultPolicy:
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 MakeQDramTileDistribution ()
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>
static CK_TILE_HOST_DEVICE constexpr auto MakeQLdsBlockDescriptor ()
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>
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::BlockFmhaFwdSplitKVPipelineNWarpSShuffleQRKSVSDefaultPolicy::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

◆ GetAlignmentOacc()

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

◆ GetAlignmentQ()

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

◆ GetSmemKPackQ()

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

◆ GetSmemNPackS()

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

◆ GetSmemSize()

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

◆ GetSmemSizeK()

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

◆ GetSmemSizeQ()

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

◆ GetSmemSizeS()

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

◆ GetSmemSizeV()

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

◆ MakeQDramTileDistribution()

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

◆ MakeQLdsBlockDescriptor()

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

◆ MakeQRegTileDistribution()

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

◆ MakeSLdsBlockDescriptor()

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

◆ MakeSRegTileDistribution()

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