Intrawave > Struct Reference

• ck_tile
• GemmPipelineAgBgCrCompAsync
• PipelineImpl< GemmPipelineScheduler::Intrawave >

Public Types | Public Member Functions | Static Public Member Functions | List of all members

ck_tile::GemmPipelineAgBgCrCompAsync< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > Struct Reference

#include <gemm_pipeline_ag_bg_cr_comp_async.hpp>

Inheritance diagram for ck_tile::GemmPipelineAgBgCrCompAsync< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave >:

Public Types
using	Base = PipelineImplBase
Public Types inherited from ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >
using	AsDataType = remove_cvref_t<typename Problem::AsDataTypeTuple>
using	BsDataType = remove_cvref_t<typename Problem::BsDataTypeTuple>
using	AsLayout = remove_cvref_t<typename Problem::AsLayoutTuple>
using	BsLayout = remove_cvref_t<typename Problem::BsLayoutTuple>
using	BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>
using	ADataType = remove_cvref_t<std::tuple_element_t<number<0>{}, AsDataType>>
using	ALayout = remove_cvref_t<std::tuple_element_t<number<0>{}, AsLayout>>
using	BDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, BsDataType>>
using	BLayout = remove_cvref_t<std::tuple_element_t<number<0>{}, BsLayout>>

Public Member Functions
template<bool HasHotLoop, TailNumber TailNum, 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 *__restrict__ p_smem_0, void *__restrict__ p_smem_1) const
Public Member Functions inherited from ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >
template<typename DstBlockTile, typename SrcTileWindow, typename DramTileWindowStep>
CK_TILE_DEVICE void	GlobalPrefetch (DstBlockTile &dst_block_tile, SrcTileWindow &dram_tile_window, const DramTileWindowStep &dram_tile_window_step) const
template<typename DstBlockWindow, typename SrcTileWindow, typename DramTileWindowStep>
CK_TILE_DEVICE void	GlobalPrefetchAsync (DstBlockWindow &dst_block_window, SrcTileWindow &dram_tile_window, const DramTileWindowStep &dram_tile_window_step) const
template<typename DstTileWindow, typename SrcBlockTile, typename ElementFunction>
CK_TILE_DEVICE void	LocalPrefill (DstTileWindow &lds_tile_window, const SrcBlockTile &src_block_tile, const ElementFunction &element_func) const
template<typename DstTileWindow, typename SrcBlockTile>
CK_TILE_DEVICE void	LocalPrefill (DstTileWindow &lds_tile_window, const SrcBlockTile &src_block_tile) const
template<typename DstBlockTile, typename SrcTileWindow, bool LoadTranspose = false>
CK_TILE_DEVICE void	LocalPrefetch (DstBlockTile &dst_block_tile, const SrcTileWindow &lds_tile_window, bool_constant< LoadTranspose >={}) const
CK_TILE_DEVICE auto	GetABLdsTensorViews (void *p_smem) const
template<typename DramBlockWindowTmp, typename std::enable_if_t< is_detected< is_tuple, DramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE constexpr auto	CopyADramWindow (const DramBlockWindowTmp &dram_block_window_tmp, const array< index_t, 2 > &offset={0, 0}) const
template<typename DramBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, DramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE constexpr auto	CopyADramWindow (const DramBlockWindowTmp &dram_block_window_tmp, const array< index_t, 2 > &offset={0, 0}) const
template<typename DramBlockWindowTmp, typename std::enable_if_t< is_detected< is_tuple, DramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE constexpr auto	CopyBDramWindow (const DramBlockWindowTmp &dram_block_window_tmp, const array< index_t, 2 > &offset={0, 0}) const
template<typename DramBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, DramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE constexpr auto	CopyBDramWindow (const DramBlockWindowTmp &dram_block_window_tmp, const array< index_t, 2 > &offset={0, 0}) const
template<typename ADramBlockWindowTmp, typename ALdsTensorView, typename ALdsLoadTileDistr>
CK_TILE_DEVICE constexpr auto	GetAWindows (const ADramBlockWindowTmp &a_dram_block_window_tmp, const ALdsTensorView &a_lds_block_view, const ALdsLoadTileDistr &, const array< index_t, 2 > &offset={0, 0}) const
template<typename BDramBlockWindowTmp, typename BLdsTensorView, typename BLdsLoadTileDistr>
CK_TILE_DEVICE constexpr auto	GetBWindows (const BDramBlockWindowTmp &b_dram_block_window_tmp, const BLdsTensorView &b_lds_block_view, const BLdsLoadTileDistr &, const array< index_t, 2 > &offset={0, 0}) const

Static Public Member Functions
static CK_TILE_DEVICE constexpr auto	HotLoopScheduler ()
Static Public Member Functions inherited from ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >
static CK_TILE_HOST_DEVICE constexpr auto	TransposeC ()

Additional Inherited Members
Static Public Attributes inherited from ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >
static constexpr index_t	MPerBlock = BlockGemmShape::kM
static constexpr index_t	NPerBlock = BlockGemmShape::kN
static constexpr index_t	KPerBlock = BlockGemmShape::kK
static constexpr bool	is_a_load_tr = false
static constexpr bool	is_b_load_tr = false

Member Typedef Documentation

◆ Base

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompAsyncDefaultPolicy>

using ck_tile::GemmPipelineAgBgCrCompAsync< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave >::Base = PipelineImplBase

Member Function Documentation

◆ HotLoopScheduler()

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompAsyncDefaultPolicy>

CK_TILE_DEVICE constexpr auto ck_tile::GemmPipelineAgBgCrCompAsync< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave >::HotLoopScheduler ( )

inlinestaticconstexpr

◆ operator()()

template<typename Problem, typename Policy = GemmPipelineAgBgCrCompAsyncDefaultPolicy>

template<bool HasHotLoop, TailNumber TailNum, 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::GemmPipelineAgBgCrCompAsync< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave >::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 *__restrict__ p_smem_0, void *__restrict__ p_smem_1 ) const

inline

---

The documentation for this struct was generated from the following file:

• gemm_pipeline_ag_bg_cr_comp_async.hpp