FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ > Struct Template Reference

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

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

#include <flatmm_kernel.hpp>

Classes

struct  FlatmmKernelArgs
 
struct  SplitKBatchOffset
 

Public Types

using TilePartitioner = remove_cvref_t< TilePartitioner_ >
 
using FlatmmPipeline = remove_cvref_t< FlatmmPipeline_ >
 
using BlockGemmShape = remove_cvref_t< typename FlatmmPipeline::BlockGemmShape >
 
using EpiloguePipeline = remove_cvref_t< EpiloguePipeline_ >
 
using ALayout = remove_cvref_t< typename FlatmmPipeline::ALayout >
 
using BLayout = remove_cvref_t< typename FlatmmPipeline::BLayout >
 
using CLayout = remove_cvref_t< typename FlatmmPipeline::CLayout >
 
using ADataType = remove_cvref_t< typename FlatmmPipeline::ADataType >
 
using BDataType = remove_cvref_t< typename FlatmmPipeline::BDataType >
 
using CDataType = remove_cvref_t< typename EpiloguePipeline::ODataType >
 

Public Member Functions

CK_TILE_DEVICE void operator() (FlatmmKernelArgs kargs) const
 

Static Public Member Functions

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 FlatmmKernelArgs MakeKernelArgs (const FlatmmHostArgs &hostArgs)
 
static constexpr CK_TILE_HOST_DEVICE index_t GetSmemSize ()
 
static CK_TILE_HOST bool IsSupportedArgument (const FlatmmKernelArgs &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, CDataType *c_ptr, const FlatmmKernelArgs &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)
 
static CK_TILE_DEVICE void RunFlatmm (const ADataType *a_ptr, const BDataType *b_flat_ptr, CDataType *c_ptr, void *smem_ptr, const FlatmmKernelArgs &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 KernelBlockSize = FlatmmPipeline::BlockSize
 
static constexpr auto I0 = number<0>()
 
static constexpr auto I1 = number<1>()
 
static constexpr auto I2 = number<2>()
 
static constexpr auto idxM = I0
 
static constexpr auto idxN = I1
 
static constexpr auto idxK = I2
 

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 = remove_cvref_t<typename FlatmmPipeline::BlockGemmShape>

◆ CDataType

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

◆ CLayout

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
using ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::CLayout = 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_ >
static constexpr CK_TILE_HOST auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::BlockSize ( )
inlinestaticconstexpr

◆ GetName()

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

◆ GetSmemSize()

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
static constexpr CK_TILE_HOST_DEVICE index_t ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::GetSmemSize ( )
inlinestaticconstexpr

◆ GridSize()

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
static constexpr CK_TILE_HOST 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_ >
static CK_TILE_HOST bool ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::IsSupportedArgument ( const FlatmmKernelArgs &  kargs)
inlinestatic

◆ MakeGemmPadViews()

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
template<typename TensorView >
static 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>
static CK_TILE_DEVICE auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::MakeGemmTensorViews ( const ADataType *  a_ptr,
const BDataType *  b_flat_ptr,
CDataType *  c_ptr,
const FlatmmKernelArgs &  kargs,
const SplitKBatchOffset &  splitk_batch_offset 
)
inlinestatic

◆ MakeGemmTileWindows()

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
template<typename PadView >
static 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_ >
static constexpr CK_TILE_HOST FlatmmKernelArgs ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::MakeKernelArgs ( const FlatmmHostArgs &  hostArgs)
inlinestaticconstexpr

◆ operator()()

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
CK_TILE_DEVICE void ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::operator() ( FlatmmKernelArgs  kargs) const
inline

◆ RunFlatmm()

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
static CK_TILE_DEVICE void ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::RunFlatmm ( const ADataType *  a_ptr,
const BDataType *  b_flat_ptr,
CDataType *  c_ptr,
void *  smem_ptr,
const FlatmmKernelArgs &  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_ >
constexpr auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::I0 = number<0>()
staticconstexpr

◆ I1

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

◆ I2

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

◆ idxK

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
constexpr auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::idxK = I2
staticconstexpr

◆ idxM

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
constexpr auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::idxM = I0
staticconstexpr

◆ idxN

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
constexpr auto ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::idxN = I1
staticconstexpr

◆ KernelBlockSize

template<typename TilePartitioner_ , typename FlatmmPipeline_ , typename EpiloguePipeline_ >
constexpr index_t ck_tile::FlatmmKernel< TilePartitioner_, FlatmmPipeline_, EpiloguePipeline_ >::KernelBlockSize = FlatmmPipeline::BlockSize
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/docs-7.0.0/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp