EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe > Struct Template Reference

EpilogueDirectStore&lt; DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe &gt; Struct Template Reference#

Composable Kernel: ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe > Struct Template Reference
Static Public Member Functions | Static Public Attributes | List of all members
ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe > Struct Template Reference

#include <epilogue_direct_store.hpp>

Static Public Member Functions

static constexpr __device__ bool IsLDSNeeded ()
 
template<InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename CThreadBuf , typename DsGridPointer , typename DsGridDesc_MBlock_MPerBlock_NBlock_NPerBlock , typename EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock >
static __device__ void Run (CThreadBuf &c_thread_buf, DsGridPointer, EDataType *p_e_grid, void *, const DsGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &, const EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &e_grid_desc_mblock_mperblock_nblock_nperblock, CDEElementwiseOperation &cde_element_op, const index_t &block_m_id, const index_t &block_n_id)
 

Static Public Attributes

static constexpr auto I0 = Number<0>{}
 
static constexpr auto I1 = Number<1>{}
 
static constexpr auto I2 = Number<2>{}
 
static constexpr auto I3 = Number<3>{}
 
static constexpr auto I4 = Number<4>{}
 
static constexpr auto I5 = Number<5>{}
 
static constexpr auto I6 = Number<6>{}
 

Member Function Documentation

◆ IsLDSNeeded()

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
static constexpr __device__ bool ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::IsLDSNeeded ( )
inlinestaticconstexpr

◆ Run()

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
template<InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename CThreadBuf , typename DsGridPointer , typename DsGridDesc_MBlock_MPerBlock_NBlock_NPerBlock , typename EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock >
static __device__ void ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::Run ( CThreadBuf &  c_thread_buf,
DsGridPointer  ,
EDataType *  p_e_grid,
void *  ,
const DsGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &  ,
const EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &  e_grid_desc_mblock_mperblock_nblock_nperblock,
CDEElementwiseOperation &  cde_element_op,
const index_t &  block_m_id,
const index_t &  block_n_id 
)
inlinestatic

Member Data Documentation

◆ I0

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
constexpr auto ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::I0 = Number<0>{}
staticconstexpr

◆ I1

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
constexpr auto ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::I1 = Number<1>{}
staticconstexpr

◆ I2

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
constexpr auto ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::I2 = Number<2>{}
staticconstexpr

◆ I3

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
constexpr auto ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::I3 = Number<3>{}
staticconstexpr

◆ I4

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
constexpr auto ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::I4 = Number<4>{}
staticconstexpr

◆ I5

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
constexpr auto ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::I5 = Number<5>{}
staticconstexpr

◆ I6

template<typename DsDataType , typename EDataType , typename AccDataType , index_t MRepeat, index_t NRepeat, typename CDEElementwiseOperation , typename BlockwiseGemmPipe >
constexpr auto ck::EpilogueDirectStore< DsDataType, EDataType, AccDataType, MRepeat, NRepeat, CDEElementwiseOperation, BlockwiseGemmPipe >::I6 = Number<6>{}
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/3643/include/ck/tensor_operation/gpu/grid/epilogue_direct_store.hpp