Intrawave > Struct Reference

Intrawave > Struct Reference#

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

#include <gemm_aquant_pipeline_ag_bg_cr_v3.hpp>

Inheritance diagram for ck_tile::AQuantGemmPipelineAgBgCrCompV3< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave >:
ck_tile::GemmAQuantPipelineAgBgCrImplBase< Problem, Policy > ck_tile::GemmAQuantPipelineAgBgCrImplBase< Problem, Policy > ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy > ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >

Public Types

using Base = PipelineImplBase
Public Types inherited from ck_tile::GemmAQuantPipelineAgBgCrImplBase< Problem, Policy >
using Base = GemmPipelineAgBgCrImplBase<Problem, Policy>
using ADataType = typename Base::ADataType
using ALayout = typename Base::ALayout
using BDataType = typename Base::BDataType
using BLayout = typename Base::BLayout
using BlockGemmShape = typename Base::BlockGemmShape
using QuantGroupSize = remove_cvref_t<typename Problem::QuantGroupSize>
using AQLayout = remove_cvref_t<typename Problem::AQLayout>
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 ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename AQDramBlockWindowTmp, 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, const AQDramBlockWindowTmp &aq_dram_block_window_tmp, index_t m, index_t num_loop, void *p_smem) const
Public Member Functions inherited from ck_tile::GemmAQuantPipelineAgBgCrImplBase< Problem, Policy >
template<typename AQDramBlockWindowTmp>
CK_TILE_DEVICE constexpr auto GetAQDramLoadWindow (const AQDramBlockWindowTmp &aq_dram_block_window_tmp) 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

Additional Inherited Members

Static Public Member Functions inherited from ck_tile::GemmPipelineAgBgCrImplBase< Problem, Policy >
static CK_TILE_HOST_DEVICE constexpr auto TransposeC ()
Static Public Attributes inherited from ck_tile::GemmAQuantPipelineAgBgCrImplBase< 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 index_t KPerBlockAQ = KPerBlock / QuantGroupSize::kK
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 = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrCompV3< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave >::Base = PipelineImplBase

Member Function Documentation

◆ operator()()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
template<bool HasHotLoop, TailNumber TailNum, typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename AQDramBlockWindowTmp, typename AElementFunction, typename BElementFunction>
CK_TILE_DEVICE auto ck_tile::AQuantGemmPipelineAgBgCrCompV3< Problem, Policy >::PipelineImpl< GemmPipelineScheduler::Intrawave >::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,
const AQDramBlockWindowTmp & aq_dram_block_window_tmp,
index_t m,
index_t num_loop,
void * p_smem ) const
inline
The documentation for this struct was generated from the following file: