Intrawave, GemmTraits > Struct Template Reference

Intrawave, GemmTraits > Struct Template Reference#

Composable Kernel: ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits > Struct Template Reference
Public Types | Public Member Functions | Public Attributes | Static Public Attributes | List of all members
ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits > Struct Template Reference

#include <block_universal_gemm_as_bs_cr.hpp>

Public Types

using ALdsTile = decltype(make_static_distributed_tensor<ComputeDataType>(ALdsTileDistr))
using BLdsTile = decltype(make_static_distributed_tensor<ComputeDataType>(BLdsTileDistr))

Public Member Functions

template<typename ASmemBlockWindow, typename BSmemBlockWindow, bool ALoadTranspose = false, bool BLoadTranspose = false>
CK_TILE_DEVICE void LocalPrefetch (const ASmemBlockWindow &a_block_window, const BSmemBlockWindow &b_block_window, bool_constant< ALoadTranspose >={}, bool_constant< BLoadTranspose >={})
template<typename CBlockTensor, typename ASmemBlockWindow, typename BSmemBlockWindow, bool ALoadTranspose = false, bool BLoadTranspose = false>
CK_TILE_DEVICE void operator() (CBlockTensor &c_block_tensor, const ASmemBlockWindow &, const BSmemBlockWindow &, bool_constant< ALoadTranspose >={}, bool_constant< BLoadTranspose >={})

Public Attributes

ALdsTile a_warp_tile_
BLdsTile b_warp_tile_

Static Public Attributes

static constexpr auto ALdsTileDistr
static constexpr auto BLdsTileDistr

Member Typedef Documentation

◆ ALdsTile

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename GemmTraits>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >::ALdsTile = decltype(make_static_distributed_tensor<ComputeDataType>(ALdsTileDistr))

◆ BLdsTile

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename GemmTraits>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >::BLdsTile = decltype(make_static_distributed_tensor<ComputeDataType>(BLdsTileDistr))

Member Function Documentation

◆ LocalPrefetch()

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename GemmTraits>
template<typename ASmemBlockWindow, typename BSmemBlockWindow, bool ALoadTranspose = false, bool BLoadTranspose = false>
CK_TILE_DEVICE void ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >::LocalPrefetch ( const ASmemBlockWindow & a_block_window,
const BSmemBlockWindow & b_block_window,
bool_constant< ALoadTranspose > = {},
bool_constant< BLoadTranspose > = {} )
inline

◆ operator()()

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename GemmTraits>
template<typename CBlockTensor, typename ASmemBlockWindow, typename BSmemBlockWindow, bool ALoadTranspose = false, bool BLoadTranspose = false>
CK_TILE_DEVICE void ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >::operator() ( CBlockTensor & c_block_tensor,
const ASmemBlockWindow & ,
const BSmemBlockWindow & ,
bool_constant< ALoadTranspose > = {},
bool_constant< BLoadTranspose > = {} )
inline

Member Data Documentation

◆ a_warp_tile_

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename GemmTraits>
ALdsTile ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >::a_warp_tile_

◆ ALdsTileDistr

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename GemmTraits>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >::ALdsTileDistr
staticconstexpr
Initial value:
=
decltype(make_static_tile_distribution(MakeABlockDistributionEncode())){}
ck_tile::make_static_tile_distribution
CK_TILE_HOST_DEVICE constexpr auto make_static_tile_distribution(StaticTileDistributionEncoding_)
Definition tile_distribution.hpp:480
ck_tile::BlockUniversalGemmAsBsCr::MakeABlockDistributionEncode
static CK_TILE_DEVICE constexpr auto MakeABlockDistributionEncode()
Definition block_universal_gemm_as_bs_cr.hpp:136

◆ b_warp_tile_

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename GemmTraits>
BLdsTile ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >::b_warp_tile_

◆ BLdsTileDistr

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename GemmTraits>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >::BLdsTileDistr
staticconstexpr
Initial value:
=
decltype(make_static_tile_distribution(MakeBBlockDistributionEncode())){}
ck_tile::BlockUniversalGemmAsBsCr::MakeBBlockDistributionEncode
static CK_TILE_DEVICE constexpr auto MakeBBlockDistributionEncode()
Definition block_universal_gemm_as_bs_cr.hpp:161
The documentation for this struct was generated from the following file: