UniversalWeightPreshufflePipelineAgBgCrPolicy Struct Reference

• ck_tile
• UniversalWeightPreshufflePipelineAgBgCrPolicy

Public Types | Static Public Member Functions | List of all members

ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy Struct Reference

#include <wp_pipeline_agmem_bgmem_creg_base_policy.hpp>

Inheritance diagram for ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy:

Public Types
using	BasePolicy = UniversalGemmBasePolicy<UniversalWeightPreshufflePipelineAgBgCrPolicy>

Static Public Member Functions
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	MakeALdsBlockDescriptor ()
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	MakeADramTileDistribution ()
template<typename Problem>
static CK_TILE_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	GetBlockWeightPreshuffle ()
template<typename Problem>
static CK_TILE_HOST_DEVICE constexpr auto	GetVectorSizeC ()
	Get the vector store size for C tensor.
Static Public Member Functions inherited from ck_tile::UniversalGemmBasePolicy< UniversalWeightPreshufflePipelineAgBgCrPolicy >
static constexpr auto	getATileAccessPattern ()
static constexpr auto	getBTileAccessPattern ()
static CK_TILE_DEVICE constexpr auto	MakeALdsBlockDescriptor ()
static CK_TILE_DEVICE constexpr auto	MakeBLdsBlockDescriptor ()
	Create LDS block descriptor for B tensor.
static CK_TILE_HOST_DEVICE constexpr auto	GetGlobalVectorLoadSize ()
	Get the maximum global memory vector load size.
static CK_TILE_HOST_DEVICE constexpr auto	GetVectorSizeA ()
static CK_TILE_HOST_DEVICE constexpr auto	GetVectorSizeB ()
static CK_TILE_HOST_DEVICE constexpr auto	GetVectorSizeC ()
	Get the vector store size for C tensor.
static CK_TILE_HOST_DEVICE constexpr auto	IsTransposeC ()
static CK_TILE_HOST_DEVICE constexpr auto	MakeADramTileDistribution ()
static CK_TILE_HOST_DEVICE constexpr auto	MakeBDramTileDistribution ()
static CK_TILE_HOST_DEVICE constexpr auto	MakeShuffledARegTileDistribution ()
static CK_TILE_HOST_DEVICE constexpr auto	MakeShuffledBRegTileDistribution ()
static CK_TILE_HOST_DEVICE constexpr auto	GetSmemPackA ()
static CK_TILE_HOST_DEVICE constexpr auto	GetSmemPackB ()
static CK_TILE_DEVICE constexpr index_t	GetSmemSizeA ()
static CK_TILE_DEVICE constexpr index_t	GetSmemSizeB ()
static CK_TILE_DEVICE constexpr index_t	GetSmemSize ()

Additional Inherited Members
Static Public Attributes inherited from ck_tile::UniversalGemmBasePolicy< UniversalWeightPreshufflePipelineAgBgCrPolicy >
static constexpr bool	is_a_load_tr
static constexpr bool	is_b_load_tr
static constexpr auto	I0
static constexpr auto	I1
static constexpr auto	I2
static constexpr auto	DefaultATileAccessPattern
static constexpr auto	DefaultBTileAccessPattern

Member Typedef Documentation

◆ BasePolicy

using ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::BasePolicy = UniversalGemmBasePolicy<UniversalWeightPreshufflePipelineAgBgCrPolicy>

Member Function Documentation

◆ GetBlockWeightPreshuffle()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::GetBlockWeightPreshuffle ( )

inlinestaticconstexpr

◆ GetKBPerLoad()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::GetKBPerLoad ( )

inlinestaticconstexpr

◆ GetSmemPackA()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::GetSmemPackA ( )

inlinestaticconstexpr

◆ GetSmemSize()

template<typename Problem>

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

inlinestaticconstexpr

◆ GetSmemSizeA()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::GetSmemSizeA ( )

inlinestaticconstexpr

◆ GetVectorSizeC()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::GetVectorSizeC ( )

inlinestaticconstexpr

Get the vector store size for C tensor.

Template Parameters

Problem

- Gemm pipeline problem class.

Note

The vector store size for output C tensor would depend on multiple factors like its data layout and warp gemm C transposition. In general it would be the number of consecutive elements in contiguous C dimension hold by single thread.

Returns

The vector store size for C tensor.

◆ MakeADramTileDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::MakeADramTileDistribution ( )

inlinestaticconstexpr

◆ MakeALdsBlockDescriptor()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::MakeALdsBlockDescriptor ( )

inlinestaticconstexpr

◆ MakeBFlatDramTileDistribution()

template<typename Problem>

CK_TILE_DEVICE constexpr auto ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::MakeBFlatDramTileDistribution ( )

inlinestaticconstexpr

◆ MakeShuffledARegBlockDistribution()

template<typename Problem>

CK_TILE_HOST_DEVICE constexpr auto ck_tile::UniversalWeightPreshufflePipelineAgBgCrPolicy::MakeShuffledARegBlockDistribution ( )

inlinestaticconstexpr

---

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

• wp_pipeline_agmem_bgmem_creg_base_policy.hpp