UniversalFlatmmPipelineAgBgCrPolicy Struct Reference

UniversalFlatmmPipelineAgBgCrPolicy Struct Reference#

Composable Kernel: ck_tile::UniversalFlatmmPipelineAgBgCrPolicy Struct Reference
ck_tile::UniversalFlatmmPipelineAgBgCrPolicy Struct Reference

#include <flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp>

Inheritance diagram for ck_tile::UniversalFlatmmPipelineAgBgCrPolicy:
ck_tile::F16xMXF4FlatmmPipelineAgBgCrPolicy ck_tile::MXF4FlatmmPipelineAgBgCrPolicy

Static Public Member Functions

template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeALdsBlockDescriptor ()
template<typename Problem, typename DataType, index_t MNPerBlock, index_t XPerTile>
static CK_TILE_HOST_DEVICE constexpr auto GetGlobalVectorLoadSize ()
 Get the maximum global memory vector load size.
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetVectorSizeA ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetVectorSizeB ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSizeA ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetSmemPackA ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetKBPerLoad ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeALDS_WarpTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeADramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeADramDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeBFlatDramTileDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto MakeShuffledARegBlockDistribution ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto GetBlockFlatmm ()

Static Public Attributes

static constexpr auto I0 = number<0>{}
static constexpr auto I1 = number<1>{}
static constexpr auto I2 = number<2>{}

Member Function Documentation

◆ GetBlockFlatmm()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::GetBlockFlatmm ( )
inlinestaticconstexpr

◆ GetGlobalVectorLoadSize()

template<typename Problem, typename DataType, index_t MNPerBlock, index_t XPerTile>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::GetGlobalVectorLoadSize ( )
inlinestaticconstexpr

Get the maximum global memory vector load size.

Template Parameters
ProblemThe UniversalGemmPipelineProblem object.
DataTypeThe tensor data type we're considering.
MNPerBlockThe MPerBlock or NPerBlock value depending on tensor (A/B).
XPerTileThe contiguous Tile dimension size.
Returns
Maximum DRAM vector load size.

◆ GetKBPerLoad()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::GetKBPerLoad ( )
inlinestaticconstexpr

◆ GetSmemPackA()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::GetSmemPackA ( )
inlinestaticconstexpr

◆ GetSmemSize()

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

◆ GetSmemSizeA()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::GetSmemSizeA ( )
inlinestaticconstexpr

◆ GetVectorSizeA()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::GetVectorSizeA ( )
inlinestaticconstexpr

◆ GetVectorSizeB()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::GetVectorSizeB ( )
inlinestaticconstexpr

◆ MakeADramDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::MakeADramDistribution ( )
inlinestaticconstexpr

◆ MakeADramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::MakeADramTileDistribution ( )
inlinestaticconstexpr

◆ MakeALDS_WarpTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::MakeALDS_WarpTileDistribution ( )
inlinestaticconstexpr

◆ MakeALdsBlockDescriptor()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::MakeALdsBlockDescriptor ( )
inlinestaticconstexpr

◆ MakeBFlatDramTileDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::MakeBFlatDramTileDistribution ( )
inlinestaticconstexpr

◆ MakeShuffledARegBlockDistribution()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::MakeShuffledARegBlockDistribution ( )
inlinestaticconstexpr

Member Data Documentation

◆ I0

auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::I0 = number<0>{}
staticconstexpr

◆ I1

auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::I1 = number<1>{}
staticconstexpr

◆ I2

auto ck_tile::UniversalFlatmmPipelineAgBgCrPolicy::I2 = number<2>{}
staticconstexpr

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