GemmPipelineAGmemBGmemCRegV2< Problem, Policy > Struct Template Reference

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

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

#include <gemm_pipeline_agmem_bgmem_creg_v2.hpp>

Public Types

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>>

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_HOST_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 ADramBlockWindowTmp, typename BDramBlockWindowTmp>
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
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 auto TransposeC ()
static CK_TILE_HOST_DEVICE constexpr index_t GetStaticLdsSize ()
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()

Static Public Attributes

static constexpr index_t APackedSize
static constexpr index_t BPackedSize
static constexpr index_t BlockSize = Problem::kBlockSize
static constexpr index_t kMPerBlock = BlockGemmShape::kM
static constexpr index_t kNPerBlock = BlockGemmShape::kN
static constexpr index_t kKPerBlock = BlockGemmShape::kK
static constexpr bool kPadM = Problem::kPadM
static constexpr bool kPadN = Problem::kPadN
static constexpr bool kPadK = Problem::kPadK
static constexpr bool Preshuffle = Problem::Preshuffle
static constexpr index_t NumWaveGroups = Problem::NumWaveGroups
static constexpr bool DoubleSmemBuffer = false

Member Typedef Documentation

◆ ADataType

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

◆ AElementWise

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::AElementWise = remove_cvref_t<typename Problem::AElementWise>

◆ ALayout

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

◆ AsDataType

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::AsDataType = remove_cvref_t<typename Problem::AsDataTypeTuple>

◆ AsLayout

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::AsLayout = remove_cvref_t<typename Problem::AsLayoutTuple>

◆ BDataType

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

◆ BElementWise

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::BElementWise = remove_cvref_t<typename Problem::BElementWise>

◆ BLayout

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

◆ BlockGemmShape

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

◆ BsDataType

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::BsDataType = remove_cvref_t<typename Problem::BsDataTypeTuple>

◆ BsLayout

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::BsLayout = remove_cvref_t<typename Problem::BsLayoutTuple>

◆ CDataType

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::CDataType = remove_cvref_t<typename Problem::CDataType>

◆ CLayout

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
using ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::CLayout = remove_cvref_t<typename Problem::CLayout>

Member Function Documentation

◆ GetName()

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

◆ GetSmemPackA()

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

◆ GetSmemPackB()

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

◆ GetSmemSize()

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

◆ GetStaticLdsSize()

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::GetStaticLdsSize ( )
inlinestaticconstexpr

◆ GetVectorSizeA()

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

◆ GetVectorSizeB()

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

◆ GetVectorSizeC()

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

◆ operator()() [1/3]

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
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::GemmPipelineAGmemBGmemCRegV2< 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()() [2/3]

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp>
CK_TILE_DEVICE auto ck_tile::GemmPipelineAGmemBGmemCRegV2< 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()() [3/3]

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
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_HOST_DEVICE auto ck_tile::GemmPipelineAGmemBGmemCRegV2< 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

◆ TransposeC()

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::TransposeC ( )
inlinestaticconstexpr

Member Data Documentation

◆ APackedSize

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
index_t ck_tile::GemmPipelineAGmemBGmemCRegV2< 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 = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
index_t ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::BlockSize = Problem::kBlockSize
staticconstexpr

◆ BPackedSize

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

◆ DoubleSmemBuffer

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
bool ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::DoubleSmemBuffer = false
staticconstexpr

◆ kKPerBlock

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
index_t ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::kKPerBlock = BlockGemmShape::kK
staticconstexpr

◆ kMPerBlock

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
index_t ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::kMPerBlock = BlockGemmShape::kM
staticconstexpr

◆ kNPerBlock

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
index_t ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::kNPerBlock = BlockGemmShape::kN
staticconstexpr

◆ kPadK

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
bool ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::kPadK = Problem::kPadK
staticconstexpr

◆ kPadM

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
bool ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::kPadM = Problem::kPadM
staticconstexpr

◆ kPadN

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
bool ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::kPadN = Problem::kPadN
staticconstexpr

◆ NumWaveGroups

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
index_t ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::NumWaveGroups = Problem::NumWaveGroups
staticconstexpr

◆ Preshuffle

template<typename Problem, typename Policy = GemmPipelineAGmemBGmemCRegV2DefaultPolicy>
bool ck_tile::GemmPipelineAGmemBGmemCRegV2< Problem, Policy >::Preshuffle = Problem::Preshuffle
staticconstexpr
The documentation for this struct was generated from the following file: