DefaultGemm2DEpilogue< Problem_, Policy_ > Struct Template Reference

DefaultGemm2DEpilogue&lt; Problem_, Policy_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ > Struct Template Reference

#include <default_2d_epilogue.hpp>

Inheritance diagram for ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >:
ck_tile::Default2DEpilogue< Problem_, void >

Public Types

using Problem = remove_cvref_t<Problem_>
using AsDataType = remove_cvref_t<typename Problem::AsDataType>
using BsDataType = remove_cvref_t<typename Problem::BsDataType>
using AccDataType = remove_cvref_t<typename Problem::AccDataType>
using ODataType = remove_cvref_t<typename Problem::ODataType>
using AsDataTypeTuple
using BsDataTypeTuple
using ADataType = remove_cvref_t<std::tuple_element_t<number<0>{}, AsDataTypeTuple>>
using BDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, BsDataTypeTuple>>
using BTypeToUse
using DsDataType = remove_cvref_t<typename Problem::DsDataType>
using DsLayout = remove_cvref_t<typename Problem::DsLayout>
using CDElementwise = remove_cvref_t<typename Problem::CDElementwise>
using CLayout = remove_cvref_t<typename Problem::CLayout>
using WG
using CWarpDstr = typename WG::CWarpDstr
Public Types inherited from ck_tile::Default2DEpilogue< Problem_, void >
using Problem
using AccDataType
using ODataType

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr auto GetVectorSizeC ()
template<index_t I>
static CK_TILE_HOST_DEVICE constexpr auto GetVectorSizeD (number< I > index)
Static Public Member Functions inherited from ck_tile::Default2DEpilogue< Problem_, void >
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()

Static Public Attributes

static constexpr bool ADataTypeIsTuple = is_detected<is_tuple, AsDataType>::value
static constexpr bool BDataTypeIsTuple = is_detected<is_tuple, BsDataType>::value
static constexpr index_t kMPerXdl = Problem::kMPerXdl
static constexpr index_t kNPerXdl = Problem::kNPerXdl
static constexpr index_t kKPerXdl = Problem::kKPerXdl
static constexpr index_t isCTransposed = Problem::isCTransposed
Static Public Attributes inherited from ck_tile::Default2DEpilogue< Problem_, void >
static constexpr bool kPadM
static constexpr bool kPadN
static constexpr bool UseRawStore
static constexpr memory_operation_enum MemoryOperation

Additional Inherited Members

Public Member Functions inherited from ck_tile::Default2DEpilogue< Problem_, void >
CK_TILE_DEVICE auto operator() (ODramWindowTmp &o_dram_window_tmp, const OAccTile &o_acc_tile, const DsDramWindows &ds_dram_windows, void *=nullptr) const

Member Typedef Documentation

◆ AccDataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::AccDataType = remove_cvref_t<typename Problem::AccDataType>

◆ ADataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::ADataType = remove_cvref_t<std::tuple_element_t<number<0>{}, AsDataTypeTuple>>

◆ AsDataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::AsDataType = remove_cvref_t<typename Problem::AsDataType>

◆ AsDataTypeTuple

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::AsDataTypeTuple
Initial value:
std::conditional_t<ADataTypeIsTuple,
remove_cvref_t<AsDataType>,
remove_cvref_t<tuple<AsDataType>>>
ck_tile::remove_cvref_t
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition type_traits.hpp:21
ck_tile::CShuffleEpilogue::ADataTypeIsTuple
static constexpr bool ADataTypeIsTuple
Definition cshuffle_epilogue.hpp:81

◆ BDataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::BDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, BsDataTypeTuple>>

◆ BsDataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::BsDataType = remove_cvref_t<typename Problem::BsDataType>

◆ BsDataTypeTuple

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::BsDataTypeTuple
Initial value:
std::conditional_t<BDataTypeIsTuple,
remove_cvref_t<BsDataType>,
remove_cvref_t<tuple<BsDataType>>>
ck_tile::CShuffleEpilogue::BDataTypeIsTuple
static constexpr bool BDataTypeIsTuple
Definition cshuffle_epilogue.hpp:82

◆ BTypeToUse

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::BTypeToUse
Initial value:
std::conditional_t<std::is_same_v<BDataType, pk_int4_t>, ADataType, BDataType>
ck_tile::CShuffleEpilogue::ADataType
remove_cvref_t< std::tuple_element_t< number< 0 >{}, AsDataTypeTuple > > ADataType
Definition cshuffle_epilogue.hpp:92
ck_tile::CShuffleEpilogue::BDataType
remove_cvref_t< std::tuple_element_t< number< 0 >{}, BsDataTypeTuple > > BDataType
Definition cshuffle_epilogue.hpp:93

◆ CDElementwise

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::CDElementwise = remove_cvref_t<typename Problem::CDElementwise>

◆ CLayout

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::CLayout = remove_cvref_t<typename Problem::CLayout>

◆ CWarpDstr

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::CWarpDstr = typename WG::CWarpDstr

◆ DsDataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::DsDataType = remove_cvref_t<typename Problem::DsDataType>

◆ DsLayout

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::DsLayout = remove_cvref_t<typename Problem::DsLayout>

◆ ODataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::ODataType = remove_cvref_t<typename Problem::ODataType>

◆ Problem

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::Problem = remove_cvref_t<Problem_>

◆ WG

template<typename Problem_, typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::WG
Initial value:
WarpGemmDispatcher<ADataType,
BTypeToUse,
AccDataType,
kMPerXdl,
kNPerXdl,
kKPerXdl,
isCTransposed>
ck_tile::WarpGemmDispatcher
typename impl::WarpGemmDispatcher< AType, BType, AccType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity, AttrNumAccess >::Type WarpGemmDispatcher
Definition warp_gemm_dispatcher.hpp:182
ck_tile::CShuffleEpilogue::isCTransposed
static constexpr index_t isCTransposed
Definition cshuffle_epilogue.hpp:111
ck_tile::CShuffleEpilogue::AccDataType
remove_cvref_t< typename Problem::AccDataType > AccDataType
Definition cshuffle_epilogue.hpp:76
ck_tile::CShuffleEpilogue::BTypeToUse
std::conditional_t< std::is_same_v< BDataType, pk_int4_t >, ADataType, BDataType > BTypeToUse
Definition cshuffle_epilogue.hpp:98
ck_tile::DefaultGemm2DEpilogue::kMPerXdl
static constexpr index_t kMPerXdl
Definition default_2d_epilogue.hpp:186
ck_tile::DefaultGemm2DEpilogue::kNPerXdl
static constexpr index_t kNPerXdl
Definition default_2d_epilogue.hpp:187
ck_tile::DefaultGemm2DEpilogue::kKPerXdl
static constexpr index_t kKPerXdl
Definition default_2d_epilogue.hpp:188
ck_tile::DefaultGemm2DEpilogue::ADataType
remove_cvref_t< std::tuple_element_t< number< 0 >{}, AsDataTypeTuple > > ADataType
Definition default_2d_epilogue.hpp:176

Member Function Documentation

◆ GetVectorSizeC()

template<typename Problem_, typename Policy_ = void>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::GetVectorSizeC ( )
inlinestaticconstexpr

◆ GetVectorSizeD()

template<typename Problem_, typename Policy_ = void>
template<index_t I>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::GetVectorSizeD ( number< I > index)
inlinestaticconstexpr

Member Data Documentation

◆ ADataTypeIsTuple

template<typename Problem_, typename Policy_ = void>
bool ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::ADataTypeIsTuple = is_detected<is_tuple, AsDataType>::value
staticconstexpr

◆ BDataTypeIsTuple

template<typename Problem_, typename Policy_ = void>
bool ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::BDataTypeIsTuple = is_detected<is_tuple, BsDataType>::value
staticconstexpr

◆ isCTransposed

template<typename Problem_, typename Policy_ = void>
index_t ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::isCTransposed = Problem::isCTransposed
staticconstexpr

◆ kKPerXdl

template<typename Problem_, typename Policy_ = void>
index_t ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::kKPerXdl = Problem::kKPerXdl
staticconstexpr

◆ kMPerXdl

template<typename Problem_, typename Policy_ = void>
index_t ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::kMPerXdl = Problem::kMPerXdl
staticconstexpr

◆ kNPerXdl

template<typename Problem_, typename Policy_ = void>
index_t ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::kNPerXdl = Problem::kNPerXdl
staticconstexpr
The documentation for this struct was generated from the following file: