FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ > Struct Template Reference

FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ > Struct Template Reference#

Composable Kernel: ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ > Struct Template Reference

#include <flatmm_kernel.hpp>

Inheritance diagram for ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >:

Classes
struct	SplitKBatchOffset

Public Types
using	TilePartitioner = remove_cvref_t<TilePartitioner_>
using	FlatmmPipeline = remove_cvref_t<FlatmmPipeline_>
using	BlockGemmShape
using	EpiloguePipeline = remove_cvref_t<EpiloguePipeline_>
using	ALayout = remove_cvref_t<typename FlatmmPipeline::ALayout>
using	BLayout = remove_cvref_t<typename FlatmmPipeline::BLayout>
using	ELayout = remove_cvref_t<typename FlatmmPipeline::CLayout>
using	DsLayout = remove_cvref_t<typename EpiloguePipeline::DsLayout>
using	DsDataType = remove_cvref_t<typename EpiloguePipeline::DsDataType>
using	ADataType = remove_cvref_t<typename FlatmmPipeline::ADataType>
using	BDataType = remove_cvref_t<typename FlatmmPipeline::BDataType>
using	EDataType = remove_cvref_t<typename EpiloguePipeline::ODataType>

Public Member Functions
template<class ScaleM, class ScaleN>
CK_TILE_DEVICE void	operator() (FlatmmKernelArgs< ScaleM, ScaleN, DsDataType::size()> kargs, int partition_idx=blockIdx.x) const

Static Public Member Functions
static CK_TILE_HOST const std::string	GetName ()
static CK_TILE_HOST constexpr auto	GridSize (index_t M, index_t N, index_t KBatch)
template<class ScaleM, class ScaleN>
static CK_TILE_HOST constexpr auto	GridSize (const FlatmmKernelArgs< ScaleM, ScaleN, DsDataType::size()> &kargs)
static CK_TILE_HOST constexpr auto	BlockSize ()
template<class ScaleM, class ScaleN>
static CK_TILE_HOST constexpr FlatmmKernelArgs< ScaleM, ScaleN, DsDataType::size()>	MakeKernelArgs (const ScaleFlatmmHostArgs< ScaleM, ScaleN, DsDataType::size()> &hostArgs)
static CK_TILE_HOST_DEVICE constexpr index_t	GetSmemPingSize ()
static CK_TILE_HOST_DEVICE constexpr index_t	GetSmemPongSize ()
template<class KernelArgs>
static CK_TILE_HOST bool	IsSupportedArgument (const KernelArgs &kargs)
template<memory_operation_enum DstInMemOp = memory_operation_enum::set, class KernelArgs>
static CK_TILE_DEVICE auto	MakeGemmTensorViews (const ADataType a_ptr, const BDataType b_flat_ptr, const std::array< const void , NumDTensor > &ds_ptr, EDataType e_ptr, const KernelArgs &kargs, const SplitKBatchOffset &splitk_batch_offset)
template<typename TensorView>
static CK_TILE_DEVICE auto	MakeGemmPadViews (const TensorView &views)
template<typename PadView>
static CK_TILE_DEVICE auto	MakeGemmTileWindows (const PadView &views, const index_t i_m, const index_t i_n)
template<class ScaleM, class ScaleN, bool UseDefaultScheduler = true>
static CK_TILE_DEVICE void	RunFlatmm (const ADataType a_ptr, const BDataType b_flat_ptr, const std::array< const void , NumDTensor > &ds_ptr, EDataType e_ptr, void smem_ptr_ping, void smem_ptr_pong, const FlatmmKernelArgs< ScaleM, ScaleN, DsDataType::size()> &kargs, const SplitKBatchOffset &splitk_batch_offset, const index_t block_idx_m, const index_t block_idx_n)

Static Public Attributes
static constexpr index_t	kBlockSize = FlatmmPipeline::BlockSize
static constexpr bool	UsePersistentKernel = FlatmmPipeline::UsePersistentKernel
static constexpr index_t	NumDTensor = DsDataType::size()
static constexpr auto	I0 = number<0>()
static constexpr auto	I1 = number<1>()
static constexpr auto	I2 = number<2>()
static constexpr auto	I3 = number<3>()

Member Typedef Documentation

◆ ADataType

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::ADataType = remove_cvref_t<typename FlatmmPipeline::ADataType>

◆ ALayout

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::ALayout = remove_cvref_t<typename FlatmmPipeline::ALayout>

◆ BDataType

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::BDataType = remove_cvref_t<typename FlatmmPipeline::BDataType>

◆ BLayout

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::BLayout = remove_cvref_t<typename FlatmmPipeline::BLayout>

◆ BlockGemmShape

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::BlockGemmShape

Initial value:

remove_cvref_t<typename FlatmmPipeline::BlockGemmShape>

ck_tile::remove_cvref_t

remove_cv_t< std::remove_reference_t< T > > remove_cvref_t

Definition type_traits.hpp:21

◆ DsDataType

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::DsDataType = remove_cvref_t<typename EpiloguePipeline::DsDataType>

◆ DsLayout

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::DsLayout = remove_cvref_t<typename EpiloguePipeline::DsLayout>

◆ EDataType

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::EDataType = remove_cvref_t<typename EpiloguePipeline::ODataType>

◆ ELayout

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::ELayout = remove_cvref_t<typename FlatmmPipeline::CLayout>

◆ EpiloguePipeline

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::EpiloguePipeline = remove_cvref_t<EpiloguePipeline_>

◆ FlatmmPipeline

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::FlatmmPipeline = remove_cvref_t<FlatmmPipeline_>

◆ TilePartitioner

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::TilePartitioner = remove_cvref_t<TilePartitioner_>

Member Function Documentation

◆ BlockSize()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST constexpr auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::BlockSize ( )

inlinestaticconstexpr

◆ GetName()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST const std::string ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::GetName ( )

inlinestaticnodiscard

◆ GetSmemPingSize()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::GetSmemPingSize ( )

inlinestaticconstexpr

◆ GetSmemPongSize()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST_DEVICE constexpr index_t ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::GetSmemPongSize ( )

inlinestaticconstexpr

◆ GridSize() [1/2]

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

template<class ScaleM, class ScaleN>

CK_TILE_HOST constexpr auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::GridSize ( const FlatmmKernelArgs< ScaleM, ScaleN, DsDataType::size()> & kargs )

inlinestaticconstexpr

◆ GridSize() [2/2]

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

CK_TILE_HOST constexpr auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::GridSize	(	index_t	M,
		index_t	N,
		index_t	KBatch )

inlinestaticconstexpr

◆ IsSupportedArgument()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

template<class KernelArgs>

CK_TILE_HOST bool ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::IsSupportedArgument ( const KernelArgs & kargs )

inlinestatic

◆ MakeGemmPadViews()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

template<typename TensorView>

CK_TILE_DEVICE auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::MakeGemmPadViews ( const TensorView & views )

inlinestatic

◆ MakeGemmTensorViews()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

template<memory_operation_enum DstInMemOp = memory_operation_enum::set, class KernelArgs>

CK_TILE_DEVICE auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::MakeGemmTensorViews	(	const ADataType *	a_ptr,
		const BDataType *	b_flat_ptr,
		const std::array< const void *, NumDTensor > &	ds_ptr,
		EDataType *	e_ptr,
		const KernelArgs &	kargs,
		const SplitKBatchOffset &	splitk_batch_offset )

inlinestatic

◆ MakeGemmTileWindows()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

template<typename PadView>

CK_TILE_DEVICE auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::MakeGemmTileWindows	(	const PadView &	views,
		const index_t	i_m,
		const index_t	i_n )

inlinestatic

◆ MakeKernelArgs()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

template<class ScaleM, class ScaleN>

CK_TILE_HOST constexpr FlatmmKernelArgs< ScaleM, ScaleN, DsDataType::size()> ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::MakeKernelArgs ( const ScaleFlatmmHostArgs< ScaleM, ScaleN, DsDataType::size()> & hostArgs )

inlinestaticconstexpr

◆ operator()()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

template<class ScaleM, class ScaleN>

CK_TILE_DEVICE void ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::operator()	(	FlatmmKernelArgs< ScaleM, ScaleN, DsDataType::size()>	kargs,
		int	partition_idx = blockIdx.x ) const

inline

◆ RunFlatmm()

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

template<class ScaleM, class ScaleN, bool UseDefaultScheduler = true>

CK_TILE_DEVICE void ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::RunFlatmm	(	const ADataType *	a_ptr,
		const BDataType *	b_flat_ptr,
		const std::array< const void *, NumDTensor > &	ds_ptr,
		EDataType *	e_ptr,
		void *	smem_ptr_ping,
		void *	smem_ptr_pong,
		const FlatmmKernelArgs< ScaleM, ScaleN, DsDataType::size()> &	kargs,
		const SplitKBatchOffset &	splitk_batch_offset,
		const index_t	block_idx_m,
		const index_t	block_idx_n )

inlinestatic

Member Data Documentation

◆ I0

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::I0 = number<0>()

staticconstexpr

◆ I1

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::I1 = number<1>()

staticconstexpr

◆ I2

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::I2 = number<2>()

staticconstexpr

◆ I3

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::I3 = number<3>()

staticconstexpr

◆ kBlockSize

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

index_t ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::kBlockSize = FlatmmPipeline::BlockSize

staticconstexpr

◆ NumDTensor

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

index_t ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::NumDTensor = DsDataType::size()

staticconstexpr

◆ UsePersistentKernel

template<typename TilePartitioner_, typename FlatmmPipeline_, typename EpiloguePipeline_>

bool ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::UsePersistentKernel = FlatmmPipeline::UsePersistentKernel

staticconstexpr

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

flatmm_kernel.hpp

FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ > Struct Template Reference

FlatmmKernel&lt; TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ &gt; Struct Template Reference#

Classes

Public Types

Public Member Functions

Static Public Member Functions

Static Public Attributes

Member Typedef Documentation

◆ ADataType

◆ ALayout

◆ BDataType

◆ BLayout

◆ BlockGemmShape

◆ DsDataType

◆ DsLayout

◆ EDataType

◆ ELayout

◆ EpiloguePipeline

◆ FlatmmPipeline

◆ TilePartitioner

Member Function Documentation

◆ BlockSize()

◆ GetName()

◆ GetSmemPingSize()

◆ GetSmemPongSize()

◆ GridSize() [1/2]

◆ GridSize() [2/2]

◆ IsSupportedArgument()

◆ MakeGemmPadViews()

◆ MakeGemmTensorViews()

◆ MakeGemmTileWindows()

◆ MakeKernelArgs()

◆ operator()()

◆ RunFlatmm()

Member Data Documentation

◆ I0

◆ I1

◆ I2

◆ I3

◆ kBlockSize

◆ NumDTensor

◆ UsePersistentKernel

FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ > Struct Template Reference#