Interwave > Struct Reference

Interwave > Struct Reference#

Composable Kernel: ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Interwave > Struct Reference
Public Types | Public Member Functions | List of all members
ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Interwave > Struct Reference

#include <gemm_pipeline_ag_bg_cr_mem.hpp>

Inheritance diagram for ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Interwave >:
ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >

Public Types

using Base = PipelineImplBase
 
- Public Types inherited from ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >
using ADataType = remove_cvref_t< typename Problem::ADataType >
 
using BDataType = remove_cvref_t< typename Problem::BDataType >
 
using ALayout = remove_cvref_t< typename Problem::ALayout >
 
using BLayout = remove_cvref_t< typename Problem::BLayout >
 
using BlockGemmShape = remove_cvref_t< typename Problem::BlockGemmShape >
 

Public Member Functions

template<bool HasHotLoop, TailNumber TailNum, typename ADramBlockWindowTmp , typename BDramBlockWindowTmp , typename AElementFunction , typename BElementFunction >
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const AElementFunction &a_element_func, const BDramBlockWindowTmp &b_dram_block_window_tmp, const BElementFunction &b_element_func, index_t num_loop, void *p_smem) 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 DstBlockTile , typename SrcTileWindow >
CK_TILE_DEVICE void LocalPrefetch (DstBlockTile &dst_block_tile, const SrcTileWindow &lds_tile_window) const
 
CK_TILE_DEVICE auto GetABLdsTensorViews (void *p_smem) const
 
template<typename ADramBlockWindowTmp , typename ALdsTensorView , typename ALdsLoadTileDistr >
constexpr CK_TILE_DEVICE 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 >
constexpr CK_TILE_DEVICE 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
 

Additional Inherited Members

- Static Public Member Functions inherited from ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >
static constexpr CK_TILE_HOST_DEVICE auto TransposeC ()
 
- 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
 

Member Typedef Documentation

◆ Base

template<typename Problem , typename Policy = UniversalGemmPipelineAgBgCrPolicy>
using ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Interwave >::Base = PipelineImplBase

Member Function Documentation

◆ operator()()

template<typename Problem , typename Policy = UniversalGemmPipelineAgBgCrPolicy>
template<bool HasHotLoop, TailNumber TailNum, typename ADramBlockWindowTmp , typename BDramBlockWindowTmp , typename AElementFunction , typename BElementFunction >
CK_TILE_DEVICE auto ck_tile::GemmPipelineAgBgCrMem< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Interwave >::operator() ( const ADramBlockWindowTmp &  a_dram_block_window_tmp,
const AElementFunction &  a_element_func,
const BDramBlockWindowTmp &  b_dram_block_window_tmp,
const BElementFunction &  b_element_func,
index_t  num_loop,
void *  p_smem 
) const
inline
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/gemm/pipeline/gemm_pipeline_ag_bg_cr_mem.hpp