#include <flatmm_kernel.hpp>
|
static CK_TILE_HOST const std::string | GetName () |
|
static constexpr CK_TILE_HOST auto | GridSize (index_t M, index_t N, index_t KBatch) |
|
static constexpr CK_TILE_HOST auto | BlockSize () |
|
static constexpr CK_TILE_HOST KernelArgs | MakeKernelArgs (const FlatmmHostArgs< NumDTensor > &hostArgs) |
|
static constexpr CK_TILE_HOST_DEVICE index_t | GetSmemSize () |
|
static CK_TILE_HOST bool | IsSupportedArgument (const KernelArgs &kargs) |
|
template<memory_operation_enum DstInMemOp = memory_operation_enum::set> |
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<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, const KernelArgs &kargs, const SplitKBatchOffset &splitk_batch_offset, const index_t block_idx_m, const index_t block_idx_n) |
|
◆ ADataType
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ ALayout
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ BDataType
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ BLayout
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ BlockGemmShape
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ DsDataType
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ DsLayout
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ EDataType
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ ELayout
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ EpiloguePipeline
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ FlatmmPipeline
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ KernelArgs
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ TilePartitioner
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ BlockSize()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ GetName()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ GetSmemSize()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ GridSize()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ IsSupportedArgument()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ MakeGemmPadViews()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
template<typename TensorView >
◆ MakeGemmTensorViews()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
template<memory_operation_enum DstInMemOp = memory_operation_enum::set>
◆ MakeGemmTileWindows()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
template<typename PadView >
◆ MakeKernelArgs()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ operator()()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ RunFlatmm()
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
template<bool UseDefaultScheduler = true>
static 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, |
|
|
const KernelArgs & |
kargs, |
|
|
const SplitKBatchOffset & |
splitk_batch_offset, |
|
|
const index_t |
block_idx_m, |
|
|
const index_t |
block_idx_n |
|
) |
| |
|
inlinestatic |
◆ I0
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ I1
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ I2
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ I3
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ kBlockSize
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
◆ NumDTensor
template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
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/flatmm/kernel/flatmm_kernel.hpp