DefaultGemm2DEpilogue< Problem_, Policy_ > Struct Template Reference

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

Composable Kernel: ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ > Struct Template Reference
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 ADataType = remove_cvref_t< typename Problem::ADataType >
 
using BDataType = remove_cvref_t< typename Problem::BDataType >
 
using AccDataType = remove_cvref_t< typename Problem::AccDataType >
 
using ODataType = remove_cvref_t< typename Problem::ODataType >
 
using BTypeToUse = std::conditional_t< std::is_same_v< BDataType, pk_int4_t >, ADataType, BDataType >
 
using DsDataType = ck_tile::tuple<>
 
using DsLayout = ck_tile::tuple<>
 
using CLayout = remove_cvref_t< typename Problem::CLayout >
 
using WG = WarpGemmMfmaDispatcher< ADataType, BTypeToUse, AccDataType, kMPerXdl, kNPerXdl, kKPerXdl, isCTransposed >
 
using CWarpDstr = typename WG::CWarpDstr
 
- Public Types inherited from ck_tile::Default2DEpilogue< Problem_, void >
using Problem = remove_cvref_t< Problem_ >
 
using AccDataType = remove_cvref_t< typename Problem::AccDataType >
 
using ODataType = remove_cvref_t< typename Problem::ODataType >
 

Static Public Member Functions

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

Static Public Attributes

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, void *=nullptr)
 
CK_TILE_DEVICE auto operator() (ODramWindowTmp &o_dram_window_tmp, const OAccTile &o_acc_tile, const DsDramWindows &, void *=nullptr)
 

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<typename Problem::ADataType>

◆ BDataType

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

◆ BTypeToUse

template<typename Problem_ , typename Policy_ = void>
using ck_tile::DefaultGemm2DEpilogue< Problem_, Policy_ >::BTypeToUse = std::conditional_t<std::is_same_v<BDataType, pk_int4_t>, ADataType, BDataType>

◆ 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 = ck_tile::tuple<>

◆ DsLayout

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

◆ 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 = WarpGemmMfmaDispatcher<ADataType, BTypeToUse, AccDataType, kMPerXdl, kNPerXdl, kKPerXdl, isCTransposed>

Member Function Documentation

◆ GetVectorSizeC()

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

◆ GetVectorSizeD()

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

Member Data Documentation

◆ isCTransposed

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

◆ kKPerXdl

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

◆ kMPerXdl

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

◆ kNPerXdl

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

The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-7.0.1/include/ck_tile/ops/epilogue/default_2d_epilogue.hpp