BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ > Struct Template Reference

BlockUniversalGemmAsBsCr&lt; Problem_, Policy_, UnaryOpSize_ &gt; Struct Template Reference#

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

#include <block_universal_gemm_as_bs_cr.hpp>

Classes

struct  BlockGemmImpl
struct  BlockGemmImpl< GemmPipelineScheduler::Default, GemmTraits >
struct  BlockGemmImpl< GemmPipelineScheduler::Intrawave, GemmTraits >
struct  BlockGemmImpl< GemmPipelineScheduler::Interwave, GemmTraits >

Public Types

using Traits = GemmTraits_<Problem_, Policy_>
using ADataType = remove_cvref_t<typename Traits::ADataType>
using BDataType = remove_cvref_t<typename Traits::BDataType>
using ComputeDataType = remove_cvref_t<typename Traits::ComputeDataType>
using CDataType = remove_cvref_t<typename Traits::CDataType>
using Loader = remove_cvref_t<InterleavedPKTypeLoader<ComputeDataType, UnaryOpSize_>>
using WarpGemm = remove_cvref_t<typename Traits::WarpGemm>
using AWarpDstr = typename WarpGemm::AWarpDstr
using BWarpDstr = typename WarpGemm::BWarpDstr
using CWarpDstr = typename WarpGemm::CWarpDstr
using AWarpTensor = typename WarpGemm::AWarpTensor
using BWarpTensor = typename WarpGemm::BWarpTensor
using CWarpTensor = typename WarpGemm::CWarpTensor
using I0 = number<0>
using I1 = number<1>

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 > a_load_tr={}, bool_constant< BLoadTranspose > b_load_tr={})
template<typename CBlockTensor, typename ASmemBlockWindow, typename BSmemBlockWindow, bool ALoadTranspose = false, bool BLoadTranspose = false>
CK_TILE_DEVICE void operator() (CBlockTensor &c_block_tensor, const ASmemBlockWindow &a_block_window, const BSmemBlockWindow &b_block_window, bool_constant< ALoadTranspose > a_load_tr={}, bool_constant< BLoadTranspose > b_load_tr={})
template<typename ASmemBlockWindow, typename BSmemBlockWindow, bool ALoadTranspose = false, bool BLoadTranspose = false>
CK_TILE_DEVICE auto operator() (const ASmemBlockWindow &a_block_window, const BSmemBlockWindow &b_block_window, bool_constant< ALoadTranspose > a_load_tr={}, bool_constant< BLoadTranspose > b_load_tr={})

Static Public Member Functions

static CK_TILE_DEVICE constexpr auto MakeABlockDistributionEncode ()
static CK_TILE_DEVICE constexpr auto MakeBBlockDistributionEncode ()
static CK_TILE_DEVICE constexpr auto MakeCBlockTile ()

Static Public Attributes

static constexpr index_t KIterPerWarp = Traits::KIterPerWarp
static constexpr index_t MIterPerWarp = Traits::MIterPerWarp
static constexpr index_t NIterPerWarp = Traits::NIterPerWarp
static constexpr index_t MWarp = Traits::MWarp
static constexpr index_t NWarp = Traits::NWarp
static constexpr auto Scheduler = Traits::Scheduler
static constexpr auto a_warp_y_lengths
static constexpr auto b_warp_y_lengths
static constexpr auto c_warp_y_lengths
static constexpr auto a_warp_y_index_zeros = uniform_sequence_gen_t<AWarpDstr::NDimY, 0>{}
static constexpr auto b_warp_y_index_zeros = uniform_sequence_gen_t<BWarpDstr::NDimY, 0>{}
static constexpr auto c_warp_y_index_zeros = uniform_sequence_gen_t<CWarpDstr::NDimY, 0>{}
static constexpr index_t APackedSize
static constexpr index_t BPackedSize

Member Typedef Documentation

◆ ADataType

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::ADataType = remove_cvref_t<typename Traits::ADataType>

◆ AWarpDstr

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::AWarpDstr = typename WarpGemm::AWarpDstr

◆ AWarpTensor

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::AWarpTensor = typename WarpGemm::AWarpTensor

◆ BDataType

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BDataType = remove_cvref_t<typename Traits::BDataType>

◆ BWarpDstr

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BWarpDstr = typename WarpGemm::BWarpDstr

◆ BWarpTensor

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BWarpTensor = typename WarpGemm::BWarpTensor

◆ CDataType

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::CDataType = remove_cvref_t<typename Traits::CDataType>

◆ ComputeDataType

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::ComputeDataType = remove_cvref_t<typename Traits::ComputeDataType>

◆ CWarpDstr

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::CWarpDstr = typename WarpGemm::CWarpDstr

◆ CWarpTensor

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::CWarpTensor = typename WarpGemm::CWarpTensor

◆ I0

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::I0 = number<0>

◆ I1

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::I1 = number<1>

◆ Loader

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::Loader = remove_cvref_t<InterleavedPKTypeLoader<ComputeDataType, UnaryOpSize_>>

◆ Traits

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::Traits = GemmTraits_<Problem_, Policy_>

◆ WarpGemm

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
using ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::WarpGemm = remove_cvref_t<typename Traits::WarpGemm>

Member Function Documentation

◆ LocalPrefetch()

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

◆ MakeABlockDistributionEncode()

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
CK_TILE_DEVICE constexpr auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::MakeABlockDistributionEncode ( )
inlinestaticconstexpr

◆ MakeBBlockDistributionEncode()

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
CK_TILE_DEVICE constexpr auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::MakeBBlockDistributionEncode ( )
inlinestaticconstexpr

◆ MakeCBlockTile()

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
CK_TILE_DEVICE constexpr auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::MakeCBlockTile ( )
inlinestaticconstexpr

◆ operator()() [1/2]

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

◆ operator()() [2/2]

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
template<typename ASmemBlockWindow, typename BSmemBlockWindow, bool ALoadTranspose = false, bool BLoadTranspose = false>
CK_TILE_DEVICE auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::operator() ( const ASmemBlockWindow & a_block_window,
const BSmemBlockWindow & b_block_window,
bool_constant< ALoadTranspose > a_load_tr = {},
bool_constant< BLoadTranspose > b_load_tr = {} )
inline

Member Data Documentation

◆ a_warp_y_index_zeros

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::a_warp_y_index_zeros = uniform_sequence_gen_t<AWarpDstr::NDimY, 0>{}
staticconstexpr

◆ a_warp_y_lengths

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::a_warp_y_lengths
staticconstexpr
Initial value:
=
to_sequence(AWarpDstr{}.get_ys_to_d_descriptor().get_lengths())
ck_tile::to_sequence
CK_TILE_HOST_DEVICE constexpr auto to_sequence(tuple< number< Is >... >)
Definition tile/core/container/sequence.hpp:1055
ck_tile::BlockUniversalGemmAsBsCr::AWarpDstr
typename WarpGemm::AWarpDstr AWarpDstr
Definition block_universal_gemm_as_bs_cr.hpp:109

◆ APackedSize

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
index_t ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::APackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<ADataType>>::PackedSize
ck_tile::numeric_traits
Definition tile/core/numeric/numeric.hpp:81

◆ b_warp_y_index_zeros

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::b_warp_y_index_zeros = uniform_sequence_gen_t<BWarpDstr::NDimY, 0>{}
staticconstexpr

◆ b_warp_y_lengths

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::b_warp_y_lengths
staticconstexpr
Initial value:
=
to_sequence(BWarpDstr{}.get_ys_to_d_descriptor().get_lengths())
ck_tile::BlockUniversalGemmAsBsCr::BWarpDstr
typename WarpGemm::BWarpDstr BWarpDstr
Definition block_universal_gemm_as_bs_cr.hpp:110

◆ BPackedSize

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
index_t ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::BPackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<BDataType>>::PackedSize

◆ c_warp_y_index_zeros

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::c_warp_y_index_zeros = uniform_sequence_gen_t<CWarpDstr::NDimY, 0>{}
staticconstexpr

◆ c_warp_y_lengths

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::c_warp_y_lengths
staticconstexpr
Initial value:
=
to_sequence(CWarpDstr{}.get_ys_to_d_descriptor().get_lengths())
ck_tile::BlockUniversalGemmAsBsCr::CWarpDstr
typename WarpGemm::CWarpDstr CWarpDstr
Definition block_universal_gemm_as_bs_cr.hpp:111

◆ KIterPerWarp

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
index_t ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::KIterPerWarp = Traits::KIterPerWarp
staticconstexpr

◆ MIterPerWarp

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
index_t ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::MIterPerWarp = Traits::MIterPerWarp
staticconstexpr

◆ MWarp

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
index_t ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::MWarp = Traits::MWarp
staticconstexpr

◆ NIterPerWarp

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
index_t ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::NIterPerWarp = Traits::NIterPerWarp
staticconstexpr

◆ NWarp

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
index_t ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::NWarp = Traits::NWarp
staticconstexpr

◆ Scheduler

template<typename Problem_, typename Policy_ = BlockGemmASmemBSmemCRegV1DefaultPolicy, index_t UnaryOpSize_ = 8>
auto ck_tile::BlockUniversalGemmAsBsCr< Problem_, Policy_, UnaryOpSize_ >::Scheduler = Traits::Scheduler
staticconstexpr
The documentation for this struct was generated from the following file: