GemmPipelineAgBgCrCompV5< Problem, Policy > Struct Template Reference

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

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

#include <gemm_pipeline_ag_bg_cr_comp_v5.hpp>

Inheritance diagram for ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >:
ck_tile::BaseGemmPipelineAgBgCrCompV5< Problem >

Classes

struct  PipelineImpl
struct  PipelineImpl< GemmPipelineScheduler::Intrawave >

Public Types

using Base = BaseGemmPipelineAgBgCrCompV5<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 ComputeDataType = remove_cvref_t<typename Problem::ComputeDataType>
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>())>
using I0 = number<0>
using I1 = number<1>
using I2 = number<2>

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_0) 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_0) 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_0) 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, const index_t num_loop, void *__restrict__ p_smem_0) 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 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::BaseGemmPipelineAgBgCrCompV5< Problem >
static CK_TILE_HOST_DEVICE constexpr auto TransposeC ()
static CK_TILE_HOST_DEVICE constexpr bool BlockHasHotloop (index_t)
static CK_TILE_HOST_DEVICE constexpr TailNumber GetBlockLoopTailNum (index_t)
template<typename RunFunction>
static CK_TILE_HOST_DEVICE auto TailHandler (const RunFunction &run_func, bool, TailNumber)

Static Public Attributes

static constexpr index_t NumWaveGroups = Problem::NumWaveGroups
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 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 index_t NumWarps = BlockGemmShape::NumWarps
static constexpr index_t KTileSize = BlockGemmShape::WarpTile::at(I2{})
Static Public Attributes inherited from ck_tile::BaseGemmPipelineAgBgCrCompV5< Problem >
static constexpr index_t PrefetchStages = 1
static constexpr index_t PrefillStages = 1
static constexpr index_t GlobalBufferNum = 1

Member Typedef Documentation

◆ ADataType

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

◆ AElementWise

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::AElementWise = remove_cvref_t<typename Problem::AElementWise>

◆ ALayout

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

◆ AsDataType

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::AsDataType = remove_cvref_t<typename Problem::AsDataTypeTuple>

◆ AsLayout

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::AsLayout = remove_cvref_t<typename Problem::AsLayoutTuple>

◆ Base

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::Base = BaseGemmPipelineAgBgCrCompV5<Problem>

◆ BDataType

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

◆ BElementWise

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::BElementWise = remove_cvref_t<typename Problem::BElementWise>

◆ BLayout

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

◆ BlockGemm

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

◆ BlockGemmShape

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

◆ BsDataType

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::BsDataType = remove_cvref_t<typename Problem::BsDataTypeTuple>

◆ BsLayout

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::BsLayout = remove_cvref_t<typename Problem::BsLayoutTuple>

◆ CDataType

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::CDataType = remove_cvref_t<typename Problem::CDataType>

◆ CLayout

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::CLayout = remove_cvref_t<typename Problem::CLayout>

◆ ComputeDataType

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::ComputeDataType = remove_cvref_t<typename Problem::ComputeDataType>

◆ I0

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::I0 = number<0>

◆ I1

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::I1 = number<1>

◆ I2

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::I2 = number<2>

◆ PipelineImplBase

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
using ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::PipelineImplBase = GemmPipelineAgBgCrImplBase<Problem, Policy>

Member Function Documentation

◆ GetName()

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

◆ GetSmemSize()

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

◆ GetVectorSizeA()

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

◆ GetVectorSizeB()

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

◆ GetVectorSizeC()

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

◆ IsTransposeC()

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

◆ operator()() [1/4]

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
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::GemmPipelineAgBgCrCompV5< 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_0 ) const
inline

◆ operator()() [2/4]

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
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::GemmPipelineAgBgCrCompV5< Problem, Policy >::operator() ( const ADramBlockWindowTmp & a_dram_block_window_tmp,
const BDramBlockWindowTmp & b_dram_block_window_tmp,
const index_t num_loop,
void *__restrict__ p_smem_0 ) const
inline

◆ operator()() [3/4]

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
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::GemmPipelineAgBgCrCompV5< 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_0 ) const
inline

◆ operator()() [4/4]

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
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::GemmPipelineAgBgCrCompV5< 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_0 ) const
inline

Member Data Documentation

◆ BlockSize

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::BlockSize = Problem::kBlockSize
staticconstexpr

◆ DoubleSmemBuffer

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::DoubleSmemBuffer = Problem::DoubleSmemBuffer
staticconstexpr

◆ HasHotLoop

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::HasHotLoop = Problem::HasHotLoop
staticconstexpr

◆ kPadK

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::kPadK = Problem::kPadK
staticconstexpr

◆ kPadM

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::kPadM = Problem::kPadM
staticconstexpr

◆ kPadN

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
bool ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::kPadN = Problem::kPadN
staticconstexpr

◆ KPerBlock

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::KPerBlock = BlockGemmShape::kK
staticconstexpr

◆ KTileSize

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::KTileSize = BlockGemmShape::WarpTile::at(I2{})
staticconstexpr

◆ MPerBlock

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::MPerBlock = BlockGemmShape::kM
staticconstexpr

◆ NPerBlock

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::NPerBlock = BlockGemmShape::kN
staticconstexpr

◆ NumWarps

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::NumWarps = BlockGemmShape::NumWarps
staticconstexpr

◆ NumWaveGroups

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::NumWaveGroups = Problem::NumWaveGroups
staticconstexpr

◆ Preshuffle

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
index_t ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::Preshuffle = Problem::Preshuffle
staticconstexpr

◆ Scheduler

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompV5DefaultPolicy>
auto ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::Scheduler = Problem::Scheduler
staticconstexpr

◆ TailNum

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