GemmPipelineAgBgCrImplBase< Problem, Policy > Struct Template Reference

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

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

#include <gemm_pipeline_ag_bg_cr_base.hpp>

Inheritance diagram for ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >:
ck_tile::GemmPipelineAgBgCrCompAsync< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrCompV3< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrCompV4< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Interwave > ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmAQuantPipelineAgBgCrImplBase< Problem, Policy > ck_tile::GemmBQuantPipelineAgBgCrImplBase< Problem, Policy > ck_tile::GemmPipelineAgBgCrCompAsync< Problem, Policy >::PipelineImpl< Scheduler > ck_tile::GemmPipelineAgBgCrCompAsync< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrCompV3< Problem, Policy >::PipelineImpl< Scheduler > ck_tile::GemmPipelineAgBgCrCompV3< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrCompV4< Problem, Policy >::PipelineImpl< Scheduler > ck_tile::GemmPipelineAgBgCrCompV4< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::PipelineImpl< Scheduler > ck_tile::GemmPipelineAgBgCrCompV5< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::PipelineImpl< Scheduler > ck_tile::GemmPipelineAgBgCrCompV6< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave > ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< Scheduler > ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Interwave > ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave >

Public Types

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<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_HOST_DEVICE constexpr auto TransposeC ()

Static Public Attributes

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

◆ ADataType

template<typename Problem, typename Policy>
using ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::ADataType = remove_cvref_t<std::tuple_element_t<number<0>{}, AsDataType>>

◆ ALayout

template<typename Problem, typename Policy>
using ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::ALayout = remove_cvref_t<std::tuple_element_t<number<0>{}, AsLayout>>

◆ AsDataType

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

◆ AsLayout

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

◆ BDataType

template<typename Problem, typename Policy>
using ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::BDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, BsDataType>>

◆ BLayout

template<typename Problem, typename Policy>
using ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::BLayout = remove_cvref_t<std::tuple_element_t<number<0>{}, BsLayout>>

◆ BlockGemmShape

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

◆ BsDataType

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

◆ BsLayout

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

Member Function Documentation

◆ CopyADramWindow() [1/2]

template<typename Problem, typename Policy>
template<typename DramBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, DramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE constexpr auto ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::CopyADramWindow ( const DramBlockWindowTmp & dram_block_window_tmp,
const array< index_t, 2 > & offset = {0, 0} ) const
inlineconstexpr

◆ CopyADramWindow() [2/2]

template<typename Problem, typename Policy>
template<typename DramBlockWindowTmp, typename std::enable_if_t< is_detected< is_tuple, DramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE constexpr auto ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::CopyADramWindow ( const DramBlockWindowTmp & dram_block_window_tmp,
const array< index_t, 2 > & offset = {0, 0} ) const
inlineconstexpr

◆ CopyBDramWindow() [1/2]

template<typename Problem, typename Policy>
template<typename DramBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, DramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE constexpr auto ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::CopyBDramWindow ( const DramBlockWindowTmp & dram_block_window_tmp,
const array< index_t, 2 > & offset = {0, 0} ) const
inlineconstexpr

◆ CopyBDramWindow() [2/2]

template<typename Problem, typename Policy>
template<typename DramBlockWindowTmp, typename std::enable_if_t< is_detected< is_tuple, DramBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE constexpr auto ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::CopyBDramWindow ( const DramBlockWindowTmp & dram_block_window_tmp,
const array< index_t, 2 > & offset = {0, 0} ) const
inlineconstexpr

◆ GetABLdsTensorViews()

template<typename Problem, typename Policy>
CK_TILE_DEVICE auto ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::GetABLdsTensorViews ( void * p_smem) const
inline

◆ GetAWindows()

template<typename Problem, typename Policy>
template<typename ADramBlockWindowTmp, typename ALdsTensorView, typename ALdsLoadTileDistr>
CK_TILE_DEVICE constexpr auto ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::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
inlineconstexpr

◆ GetBWindows()

template<typename Problem, typename Policy>
template<typename BDramBlockWindowTmp, typename BLdsTensorView, typename BLdsLoadTileDistr>
CK_TILE_DEVICE constexpr auto ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::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
inlineconstexpr

◆ GlobalPrefetch()

template<typename Problem, typename Policy>
template<typename DstBlockTile, typename SrcTileWindow, typename DramTileWindowStep>
CK_TILE_DEVICE void ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::GlobalPrefetch ( DstBlockTile & dst_block_tile,
SrcTileWindow & dram_tile_window,
const DramTileWindowStep & dram_tile_window_step ) const
inline

◆ GlobalPrefetchAsync()

template<typename Problem, typename Policy>
template<typename DstBlockWindow, typename SrcTileWindow, typename DramTileWindowStep>
CK_TILE_DEVICE void ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::GlobalPrefetchAsync ( DstBlockWindow & dst_block_window,
SrcTileWindow & dram_tile_window,
const DramTileWindowStep & dram_tile_window_step ) const
inline

◆ LocalPrefetch()

template<typename Problem, typename Policy>
template<typename DstBlockTile, typename SrcTileWindow, bool LoadTranspose = false>
CK_TILE_DEVICE void ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::LocalPrefetch ( DstBlockTile & dst_block_tile,
const SrcTileWindow & lds_tile_window,
bool_constant< LoadTranspose > = {} ) const
inline

◆ LocalPrefill() [1/2]

template<typename Problem, typename Policy>
template<typename DstTileWindow, typename SrcBlockTile>
CK_TILE_DEVICE void ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::LocalPrefill ( DstTileWindow & lds_tile_window,
const SrcBlockTile & src_block_tile ) const
inline

◆ LocalPrefill() [2/2]

template<typename Problem, typename Policy>
template<typename DstTileWindow, typename SrcBlockTile, typename ElementFunction>
CK_TILE_DEVICE void ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::LocalPrefill ( DstTileWindow & lds_tile_window,
const SrcBlockTile & src_block_tile,
const ElementFunction & element_func ) const
inline

◆ TransposeC()

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

Member Data Documentation

◆ is_a_load_tr

template<typename Problem, typename Policy>
bool ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::is_a_load_tr = false
staticconstexpr

◆ is_b_load_tr

template<typename Problem, typename Policy>
bool ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::is_b_load_tr = false
staticconstexpr

◆ KPerBlock

template<typename Problem, typename Policy>
index_t ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::KPerBlock = BlockGemmShape::kK
staticconstexpr

◆ MPerBlock

template<typename Problem, typename Policy>
index_t ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::MPerBlock = BlockGemmShape::kM
staticconstexpr

◆ NPerBlock

template<typename Problem, typename Policy>
index_t ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >::NPerBlock = BlockGemmShape::kN
staticconstexpr
The documentation for this struct was generated from the following file: