BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy Struct Reference

BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy Struct Reference#

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

#include <block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_default_policy.hpp>

Inheritance diagram for ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy:
ck_tile::BlockFmhaPipelineQXKSVSCustomPolicy< true, false, -1, 2 > ck_tile::BlockFmhaPipelineQXCustomPolicy< QLoadOnce_ >

Static Public Member Functions

template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t IsPreloadWholeNextIterationK ()
template<typename Problem>
static CK_TILE_DEVICE constexpr auto GetNumKLdsBuffers ()
template<typename Problem>
static CK_TILE_DEVICE constexpr auto GetNumPrefetchV ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetNumVLdsBuffers ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeQRegTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemKPackK ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeKLdsBlockDescriptor ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeKDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeVLdsBlockDescriptor ()
template<typename Problem>
static CK_TILE_DEVICE constexpr auto MakeVDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetQKBlockGemm ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetExclusiveKLdsBytes ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t IsFirstKLdsBufferOverlapLastVLdsBuffer ()
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 GetSmemSize ()
Static Public Member Functions inherited from ck_tile::BlockFmhaPipelineQXKSVSCustomPolicy< true, false, -1, 2 >
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 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 ()

Static Public Attributes

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

Additional Inherited Members

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

Member Function Documentation

◆ GetExclusiveKLdsBytes()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::GetExclusiveKLdsBytes ( )
inlinestaticconstexpr

◆ GetNumKLdsBuffers()

template<typename Problem>
CK_TILE_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::GetNumKLdsBuffers ( )
inlinestaticconstexpr

◆ GetNumPrefetchV()

template<typename Problem>
CK_TILE_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::GetNumPrefetchV ( )
inlinestaticconstexpr

◆ GetNumVLdsBuffers()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::GetNumVLdsBuffers ( )
inlinestaticconstexpr

◆ GetQKBlockGemm()

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

◆ GetSmemKPackK()

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

◆ GetSmemSize()

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

◆ GetSmemSizeK()

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

◆ GetSmemSizeV()

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

◆ IsFirstKLdsBufferOverlapLastVLdsBuffer()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::IsFirstKLdsBufferOverlapLastVLdsBuffer ( )
inlinestaticconstexpr

◆ IsPreloadWholeNextIterationK()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::IsPreloadWholeNextIterationK ( )
inlinestaticconstexpr

◆ MakeKDramTileDistribution()

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

◆ MakeKLdsBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::MakeKLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeQRegTileDistribution()

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

◆ MakeVDramTileDistribution()

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

◆ MakeVLdsBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::MakeVLdsBlockDescriptor ( )
inlinestaticconstexpr

Member Data Documentation

◆ NumPrefetchV

index_t ck_tile::BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy::NumPrefetchV = 2
staticconstexpr
The documentation for this struct was generated from the following file: