GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference

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

Composable Kernel: ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference

#include <gemm_multi_d_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 ALayout = remove_cvref_t<typename GemmPipeline::ALayout>
 Specify the layout configurations for A, B, E and D.
using BLayout = remove_cvref_t<typename GemmPipeline::BLayout>
using CLayout = remove_cvref_t<typename GemmPipeline::CLayout>
using DsLayout = remove_cvref_t<typename EpiloguePipeline::DsLayout>
using ADataType = remove_cvref_t<typename GemmPipeline::ADataType>
 Specify the data type configurations for A, B, E and D.
using BDataType = remove_cvref_t<typename GemmPipeline::BDataType>
using EDataType = remove_cvref_t<typename EpiloguePipeline::ODataType>
using DsDataType = remove_cvref_t<typename EpiloguePipeline::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 GemmMultiDHostArgs< 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 = 1
 ALayout and ADataType are expected to be scalars, not a tuple.
static constexpr index_t NumBTensor = 1
static constexpr index_t NumDTensor = DsDataType::size()

Member Typedef Documentation

◆ ADataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
using ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::ADataType = remove_cvref_t<typename GemmPipeline::ADataType>

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

◆ ALayout

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
using ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::ALayout = remove_cvref_t<typename GemmPipeline::ALayout>

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

◆ BDataType

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
using ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BDataType = remove_cvref_t<typename GemmPipeline::BDataType>

◆ BLayout

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
using ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BLayout = remove_cvref_t<typename GemmPipeline::BLayout>

◆ CLayout

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

◆ DsDataType

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

◆ DsLayout

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

◆ EDataType

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

◆ EpiloguePipeline

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
using ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::EpiloguePipeline = remove_cvref_t<EpiloguePipeline_>

◆ GemmPipeline

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
using ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GemmPipeline = remove_cvref_t<GemmPipeline_>

◆ TilePartitioner

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
using ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::TilePartitioner = remove_cvref_t<TilePartitioner_>

◆ UniversalGemmKernel

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
using ck_tile::GemmKernelMultiD< 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::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BlockSize ( ) ->dim3
inlinestaticconstexpr

◆ GetName()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
CK_TILE_HOST auto ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GetName ( ) ->conststd::string
inlinestatic

◆ GridSize()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
CK_TILE_HOST constexpr auto ck_tile::GemmKernelMultiD< 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::GemmKernelMultiD< 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::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::MakeKernelArgs ( const GemmMultiDHostArgs< 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::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::MaxOccupancyGridSize ( const stream_config & s) ->dim3
inlinestatic

◆ operator()()

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
CK_TILE_DEVICE auto ck_tile::GemmKernelMultiD< 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::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::kBlockSize = UniversalGemmKernel::kBlockSize
staticconstexpr

◆ NumATensor

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
index_t ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::NumATensor = 1
staticconstexpr

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

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

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 and NumBTensor have always been 1; the size of D is set by the user."

◆ NumBTensor

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
index_t ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::NumBTensor = 1
staticconstexpr

◆ NumDTensor

template<typename TilePartitioner_, typename GemmPipeline_, typename EpiloguePipeline_>
index_t ck_tile::GemmKernelMultiD< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::NumDTensor = DsDataType::size()
staticconstexpr
The documentation for this struct was generated from the following file: