GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference

GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference#

Composable Kernel: ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference

#include <gemm_multi_abd_kernel.hpp>

Public Types
using	UniversalGemmKernel
	Inject the UniversalGemmKernel base class to support execution of all necessary functions.
using	TilePartitioner = remove_cvref_t<TilePartitioner_>
using	GemmPipeline = remove_cvref_t<GemmPipeline_>
using	EpiloguePipeline = remove_cvref_t<EpiloguePipeline_>
using	AsLayout = remove_cvref_t<typename GemmPipeline::AsLayout>
	Specify the layout configurations for A, B, E and D.
using	BsLayout = remove_cvref_t<typename GemmPipeline::BsLayout>
using	CLayout = remove_cvref_t<typename GemmPipeline::CLayout>
using	DsLayout = remove_cvref_t<typename EpiloguePipeline::DsLayout>
using	AsDataType = remove_cvref_t<typename GemmPipeline::AsDataType>
	Specify the data type configurations for A, B, E and D.
using	BsDataType = remove_cvref_t<typename GemmPipeline::BsDataType>
using	EDataType = remove_cvref_t<typename EpiloguePipeline::ODataType>
using	DsDataType = remove_cvref_t<typename EpiloguePipeline::DsDataType>
using	ADataType = remove_cvref_t<std::tuple_element_t<0, AsDataType>>
using	BDataType = remove_cvref_t<std::tuple_element_t<0, BsDataType>>
using	DDataType = remove_cvref_t<std::tuple_element_t<0, DsDataType>>

Public Member Functions
CK_TILE_DEVICE auto	operator() (typename UniversalGemmKernel::KernelArgs kargs) const -> void

Static Public Member Functions
static CK_TILE_HOST auto	GetName () -> const std::string
static CK_TILE_HOST constexpr auto	GridSize (index_t M, index_t N, index_t KBatch) -> dim3
static CK_TILE_HOST auto	MaxOccupancyGridSize (const stream_config &s) -> dim3
static CK_TILE_HOST constexpr auto	BlockSize () -> dim3
static CK_TILE_HOST constexpr auto	MakeKernelArgs (const GemmMultiABDHostArgs< NumATensor, NumBTensor, NumDTensor > &hostArgs) -> typename UniversalGemmKernel::KernelArgs
static CK_TILE_HOST auto	IsSupportedArgument (const typename UniversalGemmKernel::KernelArgs &kargs) -> bool

Static Public Attributes
static constexpr index_t	kBlockSize = UniversalGemmKernel::kBlockSize
static constexpr index_t	NumATensor = AsDataType::size()
	ALayout and ADataType are expected to be a tuple, not a scalar.
static constexpr index_t	NumBTensor = BsDataType::size()
static constexpr index_t	NumDTensor = DsDataType::size()

Member Typedef Documentation

◆ ADataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::ADataType = remove_cvref_t<std::tuple_element_t<0, AsDataType>>

◆ AsDataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::AsDataType = remove_cvref_t<typename GemmPipeline::AsDataType>

Specify the data type configurations for A, B, E and D.

◆ AsLayout

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::AsLayout = remove_cvref_t<typename GemmPipeline::AsLayout>

Specify the layout configurations for A, B, E and D.

◆ BDataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BDataType = remove_cvref_t<std::tuple_element_t<0, BsDataType>>

◆ BsDataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BsDataType = remove_cvref_t<typename GemmPipeline::BsDataType>

◆ BsLayout

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BsLayout = remove_cvref_t<typename GemmPipeline::BsLayout>

◆ CLayout

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::CLayout = remove_cvref_t<typename GemmPipeline::CLayout>

◆ DDataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::DDataType = remove_cvref_t<std::tuple_element_t<0, DsDataType>>

◆ DsDataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::DsDataType = remove_cvref_t<typename EpiloguePipeline::DsDataType>

◆ DsLayout

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::DsLayout = remove_cvref_t<typename EpiloguePipeline::DsLayout>

◆ EDataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::EDataType = remove_cvref_t<typename EpiloguePipeline::ODataType>

◆ EpiloguePipeline

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::EpiloguePipeline = remove_cvref_t<EpiloguePipeline_>

◆ GemmPipeline

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GemmPipeline = remove_cvref_t<GemmPipeline_>

◆ TilePartitioner

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::TilePartitioner = remove_cvref_t<TilePartitioner_>

◆ UniversalGemmKernel

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

using ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::UniversalGemmKernel

Initial value:

UniversalGemmKernel<TilePartitioner_, GemmPipeline_, EpiloguePipeline_>

ck_tile::BatchedGemmKernel::UniversalGemmKernel

UniversalGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > UniversalGemmKernel

Inject the UniversalGemmKernel base class to support execution of all necessary functions.

Definition batched_gemm_kernel.hpp:65

Inject the UniversalGemmKernel base class to support execution of all necessary functions.

Member Function Documentation

◆ BlockSize()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST constexpr auto ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BlockSize ( ) ->dim3

inlinestaticconstexpr

◆ GetName()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST auto ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GetName ( ) ->conststd::string

inlinestatic

◆ GridSize()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST constexpr auto ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GridSize	(	index_t	M,
		index_t	N,
		index_t	KBatch )->dim3

inlinestaticconstexpr

◆ IsSupportedArgument()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST auto ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::IsSupportedArgument ( const typename UniversalGemmKernel::KernelArgs & kargs ) ->bool

inlinestatic

◆ MakeKernelArgs()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST constexpr auto ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::MakeKernelArgs ( const GemmMultiABDHostArgs< NumATensor, NumBTensor, NumDTensor > & hostArgs ) ->typenameUniversalGemmKernel::KernelArgs

inlinestaticconstexpr

Universal GEMM requires array objects and corresponding stride information for matrices A, B, and D.

◆ MaxOccupancyGridSize()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST auto ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::MaxOccupancyGridSize ( const stream_config & s ) ->dim3

inlinestatic

◆ operator()()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

CK_TILE_DEVICE auto ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::operator() ( typename UniversalGemmKernel::KernelArgs kargs ) const->void

inline

Member Data Documentation

◆ kBlockSize

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

index_t ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::kBlockSize = UniversalGemmKernel::kBlockSize

staticconstexpr

◆ NumATensor

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

index_t ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::NumATensor = AsDataType::size()

staticconstexpr

ALayout and ADataType are expected to be a tuple, not a scalar.

BLayout and BDataType are expected to be a tuple, not a scalar.

CLayout and EDataType are expected to be scalars, not a tuple.

DsLayout and DsDataType are expected to be tuple, not a scalar.

The sizes of NumATensor, NumBTensor and NumDTensor is set by the user."

◆ NumBTensor

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

index_t ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::NumBTensor = BsDataType::size()

staticconstexpr

◆ NumDTensor

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>

index_t ck_tile::GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::NumDTensor = DsDataType::size()

staticconstexpr

The documentation for this struct was generated from the following file:

gemm_multi_abd_kernel.hpp

GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference

GemmKernelMultiABD&lt; TilePartitioner_, GemmPipeline_, EpiloguePipeline_ &gt; Struct Template Reference#

Public Types

Public Member Functions

Static Public Member Functions

Static Public Attributes

Member Typedef Documentation

◆ ADataType

◆ AsDataType

◆ AsLayout

◆ BDataType

◆ BsDataType

◆ BsLayout

◆ CLayout

◆ DDataType

◆ DsDataType

◆ DsLayout

◆ EDataType

◆ EpiloguePipeline

◆ GemmPipeline

◆ TilePartitioner

◆ UniversalGemmKernel

Member Function Documentation

◆ BlockSize()

◆ GetName()

◆ GridSize()

◆ IsSupportedArgument()

◆ MakeKernelArgs()

◆ MaxOccupancyGridSize()

◆ operator()()

Member Data Documentation

◆ kBlockSize

◆ NumATensor

◆ NumBTensor

◆ NumDTensor

GemmKernelMultiABD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference#