GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference#
Public Types |
Public Member Functions |
Static Public Member Functions |
Static Public Attributes |
List of all members
ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference
#include <gemm_kernel.hpp>
Public Types | |
using | UniversalGemmKernel = UniversalGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > |
Inject the UniversalGemmKernel base class to support execution of all necessary functions. More... | |
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. More... | |
using | BLayout = remove_cvref_t< typename GemmPipeline::BLayout > |
using | ELayout = remove_cvref_t< typename GemmPipeline::CLayout > |
using | ADataType = remove_cvref_t< typename GemmPipeline::ADataType > |
Specify the data type configurations for A, B, E and D. More... | |
using | BDataType = remove_cvref_t< typename GemmPipeline::BDataType > |
using | EDataType = remove_cvref_t< typename EpiloguePipeline::ODataType > |
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 constexpr CK_TILE_HOST auto | GridSize (index_t M, index_t N, index_t KBatch) -> dim3 |
static CK_TILE_HOST auto | MaxOccupancyGridSize (const stream_config &s) -> dim3 |
static constexpr CK_TILE_HOST auto | BlockSize () -> dim3 |
static constexpr CK_TILE_HOST auto | MakeKernelArgs (const GemmHostArgs &hostArgs) -> typename UniversalGemmKernel::KernelArgs |
static CK_TILE_HOST auto | IsSupportedArgument (const typename UniversalGemmKernel::KernelArgs &kargs) -> bool |
Static Public Attributes | |
static constexpr index_t | NumATensor = 1 |
ALayout and ADataType are expected to be scalars, not a tuple. More... | |
static constexpr index_t | NumBTensor = 1 |
static constexpr index_t | kBlockSize = UniversalGemmKernel::kBlockSize |
Member Typedef Documentation
◆ ADataType
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::GemmKernel< 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::GemmKernel< 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::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BDataType = remove_cvref_t<typename GemmPipeline::BDataType> |
◆ BLayout
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BLayout = remove_cvref_t<typename GemmPipeline::BLayout> |
◆ EDataType
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::EDataType = remove_cvref_t<typename EpiloguePipeline::ODataType> |
◆ ELayout
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::ELayout = remove_cvref_t<typename GemmPipeline::CLayout> |
◆ EpiloguePipeline
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::EpiloguePipeline = remove_cvref_t<EpiloguePipeline_> |
◆ GemmPipeline
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GemmPipeline = remove_cvref_t<GemmPipeline_> |
◆ TilePartitioner
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::TilePartitioner = remove_cvref_t<TilePartitioner_> |
◆ UniversalGemmKernel
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::UniversalGemmKernel = UniversalGemmKernel<TilePartitioner_, GemmPipeline_, EpiloguePipeline_> |
Inject the UniversalGemmKernel base class to support execution of all necessary functions.
Member Function Documentation
◆ BlockSize()
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
inlinestaticconstexpr |
◆ GetName()
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
inlinestatic |
◆ GridSize()
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
inlinestaticconstexpr |
◆ IsSupportedArgument()
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
inlinestatic |
◆ MakeKernelArgs()
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
inlinestaticconstexpr |
Universal GEMM requires array objects and corresponding stride information for matrices A, B.
◆ MaxOccupancyGridSize()
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
inlinestatic |
◆ operator()()
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
inline |
Member Data Documentation
◆ kBlockSize
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
staticconstexpr |
◆ NumATensor
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
staticconstexpr |
ALayout and ADataType are expected to be scalars, not a tuple.
BLayout and BDataType are expected to be scalars, not a tuple.
C/ELayout and C/EDataType are expected to be scalars, not a tuple.
◆ NumBTensor
template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >
|
staticconstexpr |
The documentation for this struct was generated from the following file:
- /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp