GemmPipelineAgBgCrCompV6< Problem, Policy > Struct Template Reference

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

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

#include <gemm_pipeline_ag_bg_cr_comp_v6.hpp>

Inheritance diagram for ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >:
ck_tile::BaseGemmPipelineAgBgCrCompV6< Problem >

Classes

struct  PipelineImpl
struct  PipelineImpl< GemmPipelineScheduler::Intrawave >

Public Types

using Base = BaseGemmPipelineAgBgCrCompV6<Problem>
using BasePImpl = 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 BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>
using AElementWise = remove_cvref_t<typename Problem::AElementWise>
using BElementWise = remove_cvref_t<typename Problem::BElementWise>
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>())>

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 *__restrict__ 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, const index_t num_loop, void *__restrict__ 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 *__restrict__ 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 CK_TILE_HOST_DEVICE constexpr auto IsTransposeC ()
Static Public Member Functions inherited from ck_tile::BaseGemmPipelineAgBgCrCompV6< Problem >
static CK_TILE_HOST_DEVICE constexpr auto TransposeC ()
static CK_TILE_HOST constexpr bool BlockHasHotloop (index_t num_loop)
static CK_TILE_HOST 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 APackedSize
static constexpr index_t BPackedSize
static constexpr index_t NumWaveGroups = Problem::NumWaveGroups
static constexpr auto I0 = number<0>{}
static constexpr auto I1 = number<1>{}
static constexpr auto I2 = number<2>{}
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 KRepeat = BlockGemm::WarpGemm::kKPerThread / GetSmemPackA()
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 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<BasePImpl::is_a_load_tr>{}
static constexpr auto is_b_load_tr_v = bool_constant<BasePImpl::is_b_load_tr>{}
Static Public Attributes inherited from ck_tile::BaseGemmPipelineAgBgCrCompV6< Problem >
static constexpr index_t PrefetchStages = 3
static constexpr index_t PrefillStages = 1
static constexpr index_t GlobalBufferNum = 2
static constexpr index_t HotloopUnroll = 2

Member Typedef Documentation

◆ ADataType

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

◆ AElementWise

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::AElementWise = remove_cvref_t<typename Problem::AElementWise>

◆ ALayout

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

◆ AsDataType

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::AsDataType = remove_cvref_t<typename Problem::AsDataTypeTuple>

◆ AsLayout

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::AsLayout = remove_cvref_t<typename Problem::AsLayoutTuple>

◆ Base

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::Base = BaseGemmPipelineAgBgCrCompV6<Problem>

◆ BasePImpl

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::BasePImpl = GemmPipelineAgBgCrImplBase<Problem, Policy>

◆ BDataType

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

◆ BElementWise

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::BElementWise = remove_cvref_t<typename Problem::BElementWise>

◆ BLayout

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

◆ BlockGemm

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

◆ BlockGemmShape

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

◆ BsDataType

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::BsDataType = remove_cvref_t<typename Problem::BsDataTypeTuple>

◆ BsLayout

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::BsLayout = remove_cvref_t<typename Problem::BsLayoutTuple>

◆ CDataType

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::CDataType = remove_cvref_t<typename Problem::CDataType>

◆ CLayout

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::CLayout = remove_cvref_t<typename Problem::CLayout>

Member Function Documentation

◆ GetName()

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

◆ GetSmemPackA()

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

◆ GetSmemPackB()

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

◆ GetSmemSize()

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

◆ GetVectorSizeA()

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

◆ GetVectorSizeB()

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

◆ GetVectorSizeC()

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

◆ IsTransposeC()

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::IsTransposeC ( )
inlinestaticconstexpr

◆ operator()() [1/3]

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
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::GemmPipelineAgBgCrCompV6< 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 *__restrict__ p_smem ) const
inline

◆ operator()() [2/3]

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
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::GemmPipelineAgBgCrCompV6< 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 *__restrict__ p_smem ) const
inline

◆ operator()() [3/3]

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
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::GemmPipelineAgBgCrCompV6< Problem, Policy >::operator() ( const AsDramBlockWindowTmp & a_dram_block_window_tmp,
const BsDramBlockWindowTmp & b_dram_block_window_tmp,
const index_t num_loop,
void *__restrict__ p_smem ) const
inline

Member Data Documentation

◆ APackedSize

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::APackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<ADataType>>::PackedSize
ck_tile::numeric_traits
Definition tile/core/numeric/numeric.hpp:81

◆ BlockSize

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::BlockSize = Problem::kBlockSize
staticconstexpr

◆ BPackedSize

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::BPackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<BDataType>>::PackedSize

◆ DoubleSmemBuffer

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::DoubleSmemBuffer = Problem::DoubleSmemBuffer
staticconstexpr

◆ HasHotLoop

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::HasHotLoop = Problem::HasHotLoop
staticconstexpr

◆ I0

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
auto ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::I0 = number<0>{}
staticconstexpr

◆ I1

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
auto ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::I1 = number<1>{}
staticconstexpr

◆ I2

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
auto ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::I2 = number<2>{}
staticconstexpr

◆ is_a_load_tr_v

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
auto ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::is_a_load_tr_v = bool_constant<BasePImpl::is_a_load_tr>{}
staticconstexpr

◆ is_b_load_tr_v

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
auto ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::is_b_load_tr_v = bool_constant<BasePImpl::is_b_load_tr>{}
staticconstexpr

◆ kPadK

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::kPadK = Problem::kPadK
staticconstexpr

◆ kPadM

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::kPadM = Problem::kPadM
staticconstexpr

◆ kPadN

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::kPadN = Problem::kPadN
staticconstexpr

◆ KPerBlock

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::KPerBlock = BlockGemmShape::kK
staticconstexpr

◆ KRepeat

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::KRepeat = BlockGemm::WarpGemm::kKPerThread / GetSmemPackA()
staticconstexpr

◆ MPerBlock

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::MPerBlock = BlockGemmShape::kM
staticconstexpr

◆ NPerBlock

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::NPerBlock = BlockGemmShape::kN
staticconstexpr

◆ NumWaveGroups

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::NumWaveGroups = Problem::NumWaveGroups
staticconstexpr

◆ Preshuffle

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::Preshuffle = Problem::Preshuffle
staticconstexpr

◆ Scheduler

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV6DefaultPolicy>
auto ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::Scheduler = Problem::Scheduler
staticconstexpr

◆ TailNum

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