GemmPipelineAgBgCrMem< Problem, Policy > Struct Template Reference

GemmPipelineAgBgCrMem&lt; Problem, Policy &gt; Struct Template Reference#

Composable Kernel: ck_tile::GemmPipelineAgBgCrMem< Problem, Policy > Struct Template Reference
Classes | Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::GemmPipelineAgBgCrMem< Problem, Policy > Struct Template Reference

#include <gemm_pipeline_ag_bg_cr_mem.hpp>

Inheritance diagram for ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >:
ck_tile::BaseGemmPipelineAgBgCrMem< Problem >

Classes

struct  PipelineImpl
struct  PipelineImpl< GemmPipelineScheduler::Intrawave >
struct  PipelineImpl< GemmPipelineScheduler::Interwave >

Public Types

using Base = BaseGemmPipelineAgBgCrMem<Problem>
using PipelineImplBase = GemmPipelineAgBgCrImplBase<Problem, Policy>
using AsDataType = remove_cvref_t<typename Problem::AsDataTypeTuple>
using BsDataType = remove_cvref_t<typename Problem::BsDataTypeTuple>
using CDataType = remove_cvref_t<typename Problem::CDataType>
using AElementWise = remove_cvref_t<typename Problem::AElementWise>
using BElementWise = remove_cvref_t<typename Problem::BElementWise>
using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>
using AsLayout = remove_cvref_t<typename Problem::AsLayoutTuple>
using BsLayout = remove_cvref_t<typename Problem::BsLayoutTuple>
using CLayout = remove_cvref_t<typename Problem::CLayout>
using ALayout = remove_cvref_t<std::tuple_element_t<0, AsLayout>>
using BLayout = remove_cvref_t<std::tuple_element_t<0, BsLayout>>
using ADataType = remove_cvref_t<std::tuple_element_t<0, AsDataType>>
using BDataType = remove_cvref_t<std::tuple_element_t<0, BsDataType>>
using BlockGemm = remove_cvref_t<decltype(Policy::template GetBlockGemm<Problem>())>
using I0 = number<0>
using I1 = number<1>
using I2 = number<2>
Public Types inherited from ck_tile::BaseGemmPipelineAgBgCrMem< Problem >
using ADataType = remove_cvref_t<typename Problem::ADataType>
using BDataType = remove_cvref_t<typename Problem::BDataType>
using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

Public Member Functions

template<typename AsDramBlockWindowTmp, typename BsDramBlockWindowTmp, typename AElementFunction, typename BElementFunction, typename std::enable_if_t< is_detected< is_tuple, AsDramBlockWindowTmp >::value &&is_detected< is_tuple, BsDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const AsDramBlockWindowTmp &a_dram_block_window_tmp, const AElementFunction &a_element_func, const BsDramBlockWindowTmp &b_dram_block_window_tmp, const BElementFunction &b_element_func, index_t num_loop, void *p_smem) const
template<typename AsDramBlockWindowTmp, typename BsDramBlockWindowTmp, typename std::enable_if_t< is_detected< is_tuple, AsDramBlockWindowTmp >::value &&is_detected< is_tuple, BsDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const AsDramBlockWindowTmp &a_dram_block_window_tmp, const BsDramBlockWindowTmp &b_dram_block_window_tmp, index_t num_loop, bool has_hot_loop, TailNumber tail_number, void *p_smem) const
template<typename AsDramBlockWindowTmp, typename BsDramBlockWindowTmp, typename std::enable_if_t< is_detected< is_tuple, AsDramBlockWindowTmp >::value &&is_detected< is_tuple, BsDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const AsDramBlockWindowTmp &a_dram_block_window_tmp, const BsDramBlockWindowTmp &b_dram_block_window_tmp, index_t num_loop, void *p_smem) const
template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename AElementFunction, typename BElementFunction, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const AElementFunction &a_element_func, const BDramBlockWindowTmp &b_dram_block_window_tmp, const BElementFunction &b_element_func, index_t num_loop, void *p_smem) const
template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const BDramBlockWindowTmp &b_dram_block_window_tmp, index_t num_loop, bool has_hot_loop, TailNumber tail_number, void *p_smem) const
template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const BDramBlockWindowTmp &b_dram_block_window_tmp, index_t num_loop, void *p_smem) const

Static Public Member Functions

template<bool IsWave32Host = false>
static constexpr index_t GetVectorSizeA ()
template<bool IsWave32Host = false>
static constexpr index_t GetVectorSizeB ()
static constexpr index_t GetVectorSizeC ()
static constexpr index_t GetSmemPackA ()
static constexpr index_t GetSmemPackB ()
static CK_TILE_HOST const std::string GetName ()
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()
Static Public Member Functions inherited from ck_tile::BaseGemmPipelineAgBgCrMem< Problem >
static CK_TILE_HOST_DEVICE constexpr auto TransposeC ()
static CK_TILE_HOST_DEVICE constexpr bool BlockHasHotloop (index_t num_loop)
static CK_TILE_HOST_DEVICE constexpr TailNumber GetBlockLoopTailNum (index_t num_loop)
template<typename RunFunction>
static CK_TILE_HOST_DEVICE auto TailHandler (const RunFunction &run_func, bool has_hot_loop, TailNumber tail_number)

Static Public Attributes

static constexpr index_t MPerBlock = BlockGemmShape::kM
static constexpr index_t NPerBlock = BlockGemmShape::kN
static constexpr index_t KPerBlock = BlockGemmShape::kK
static constexpr bool kPadM = Problem::kPadM
static constexpr bool kPadN = Problem::kPadN
static constexpr bool kPadK = Problem::kPadK
static constexpr bool DoubleSmemBuffer = Problem::DoubleSmemBuffer
static constexpr index_t NumWaveGroups = Problem::NumWaveGroups
static constexpr index_t Preshuffle = Problem::Preshuffle
static constexpr bool HasHotLoop = Problem::HasHotLoop
static constexpr auto TailNum = Problem::TailNum
static constexpr auto Scheduler = Problem::Scheduler
static constexpr auto is_a_load_tr_v = bool_constant<PipelineImplBase::is_a_load_tr>{}
static constexpr auto is_b_load_tr_v = bool_constant<PipelineImplBase::is_b_load_tr>{}
static constexpr index_t PrefetchStages
Static Public Attributes inherited from ck_tile::BaseGemmPipelineAgBgCrMem< Problem >
static constexpr index_t APackedSize
static constexpr index_t BPackedSize
static constexpr index_t BlockSize = Problem::kBlockSize
static constexpr index_t MPerBlock = BlockGemmShape::kM
static constexpr index_t NPerBlock = BlockGemmShape::kN
static constexpr index_t KPerBlock = BlockGemmShape::kK
static constexpr index_t MinMemInFlyBytes = 32768
static constexpr index_t WgpPerCU
static constexpr index_t FullMemBandPrefetchStages
static constexpr index_t PrefetchStages
static constexpr index_t LocalPrefillStages = 1
static constexpr index_t GlobalBufferNum = PrefetchStages
static constexpr bool UsePersistentKernel = Problem::Traits::UsePersistentKernel

Member Typedef Documentation

◆ ADataType

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::ADataType = remove_cvref_t<std::tuple_element_t<0, AsDataType>>

◆ AElementWise

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::AElementWise = remove_cvref_t<typename Problem::AElementWise>

◆ ALayout

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::ALayout = remove_cvref_t<std::tuple_element_t<0, AsLayout>>

◆ AsDataType

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::AsDataType = remove_cvref_t<typename Problem::AsDataTypeTuple>

◆ AsLayout

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::AsLayout = remove_cvref_t<typename Problem::AsLayoutTuple>

◆ Base

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::Base = BaseGemmPipelineAgBgCrMem<Problem>

◆ BDataType

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::BDataType = remove_cvref_t<std::tuple_element_t<0, BsDataType>>

◆ BElementWise

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::BElementWise = remove_cvref_t<typename Problem::BElementWise>

◆ BLayout

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::BLayout = remove_cvref_t<std::tuple_element_t<0, BsLayout>>

◆ BlockGemm

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::BlockGemm = remove_cvref_t<decltype(Policy::template GetBlockGemm<Problem>())>

◆ BlockGemmShape

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

◆ BsDataType

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::BsDataType = remove_cvref_t<typename Problem::BsDataTypeTuple>

◆ BsLayout

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::BsLayout = remove_cvref_t<typename Problem::BsLayoutTuple>

◆ CDataType

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::CDataType = remove_cvref_t<typename Problem::CDataType>

◆ CLayout

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::CLayout = remove_cvref_t<typename Problem::CLayout>

◆ I0

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::I0 = number<0>

◆ I1

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::I1 = number<1>

◆ I2

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::I2 = number<2>

◆ PipelineImplBase

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImplBase = GemmPipelineAgBgCrImplBase<Problem, Policy>

Member Function Documentation

◆ GetName()

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
CK_TILE_HOST const std::string ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::GetName ( )
inlinestaticnodiscard

◆ GetSmemPackA()

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
constexpr index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::GetSmemPackA ( )
inlinestaticconstexpr

◆ GetSmemPackB()

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
constexpr index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::GetSmemPackB ( )
inlinestaticconstexpr

◆ GetSmemSize()

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::GetSmemSize ( )
inlinestaticconstexpr

◆ GetVectorSizeA()

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<bool IsWave32Host = false>
constexpr index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::GetVectorSizeA ( )
inlinestaticconstexpr

◆ GetVectorSizeB()

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<bool IsWave32Host = false>
constexpr index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::GetVectorSizeB ( )
inlinestaticconstexpr

◆ GetVectorSizeC()

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
constexpr index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::GetVectorSizeC ( )
inlinestaticconstexpr

◆ operator()() [1/6]

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename AElementFunction, typename BElementFunction, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::operator() ( const ADramBlockWindowTmp & a_dram_block_window_tmp,
const AElementFunction & a_element_func,
const BDramBlockWindowTmp & b_dram_block_window_tmp,
const BElementFunction & b_element_func,
index_t num_loop,
void * p_smem ) const
inline

◆ operator()() [2/6]

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::operator() ( const ADramBlockWindowTmp & a_dram_block_window_tmp,
const BDramBlockWindowTmp & b_dram_block_window_tmp,
index_t num_loop,
bool has_hot_loop,
TailNumber tail_number,
void * p_smem ) const
inline

◆ operator()() [3/6]

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::operator() ( const ADramBlockWindowTmp & a_dram_block_window_tmp,
const BDramBlockWindowTmp & b_dram_block_window_tmp,
index_t num_loop,
void * p_smem ) const
inline

◆ operator()() [4/6]

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<typename AsDramBlockWindowTmp, typename BsDramBlockWindowTmp, typename AElementFunction, typename BElementFunction, typename std::enable_if_t< is_detected< is_tuple, AsDramBlockWindowTmp >::value &&is_detected< is_tuple, BsDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::operator() ( const AsDramBlockWindowTmp & a_dram_block_window_tmp,
const AElementFunction & a_element_func,
const BsDramBlockWindowTmp & b_dram_block_window_tmp,
const BElementFunction & b_element_func,
index_t num_loop,
void * p_smem ) const
inline

◆ operator()() [5/6]

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<typename AsDramBlockWindowTmp, typename BsDramBlockWindowTmp, typename std::enable_if_t< is_detected< is_tuple, AsDramBlockWindowTmp >::value &&is_detected< is_tuple, BsDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::operator() ( const AsDramBlockWindowTmp & a_dram_block_window_tmp,
const BsDramBlockWindowTmp & b_dram_block_window_tmp,
index_t num_loop,
bool has_hot_loop,
TailNumber tail_number,
void * p_smem ) const
inline

◆ operator()() [6/6]

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<typename AsDramBlockWindowTmp, typename BsDramBlockWindowTmp, typename std::enable_if_t< is_detected< is_tuple, AsDramBlockWindowTmp >::value &&is_detected< is_tuple, BsDramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::operator() ( const AsDramBlockWindowTmp & a_dram_block_window_tmp,
const BsDramBlockWindowTmp & b_dram_block_window_tmp,
index_t num_loop,
void * p_smem ) const
inline

Member Data Documentation

◆ DoubleSmemBuffer

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
bool ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::DoubleSmemBuffer = Problem::DoubleSmemBuffer
staticconstexpr

◆ HasHotLoop

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
bool ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::HasHotLoop = Problem::HasHotLoop
staticconstexpr

◆ is_a_load_tr_v

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::is_a_load_tr_v = bool_constant<PipelineImplBase::is_a_load_tr>{}
staticconstexpr

◆ is_b_load_tr_v

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::is_b_load_tr_v = bool_constant<PipelineImplBase::is_b_load_tr>{}
staticconstexpr

◆ kPadK

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
bool ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::kPadK = Problem::kPadK
staticconstexpr

◆ kPadM

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
bool ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::kPadM = Problem::kPadM
staticconstexpr

◆ kPadN

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
bool ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::kPadN = Problem::kPadN
staticconstexpr

◆ KPerBlock

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::KPerBlock = BlockGemmShape::kK
staticconstexpr

◆ MPerBlock

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::MPerBlock = BlockGemmShape::kM
staticconstexpr

◆ NPerBlock

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::NPerBlock = BlockGemmShape::kN
staticconstexpr

◆ NumWaveGroups

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::NumWaveGroups = Problem::NumWaveGroups
staticconstexpr

◆ PrefetchStages

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::PrefetchStages
staticconstexpr

◆ Preshuffle

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
index_t ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::Preshuffle = Problem::Preshuffle
staticconstexpr

◆ Scheduler

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::Scheduler = Problem::Scheduler
staticconstexpr

◆ TailNum

template<typename Problem, typename Policy = UniversalGemmPipelineAgBgCrPolicy>
auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::TailNum = Problem::TailNum
staticconstexpr
The documentation for this struct was generated from the following file: