FlatmmSn_32x128x512_1x4x1_16x16x32_BF16 Struct Reference

FlatmmSn_32x128x512_1x4x1_16x16x32_BF16 Struct Reference#

Composable Kernel: ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_BF16 Struct Reference
Public Types | Public Member Functions | List of all members
ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_BF16 Struct Reference

#include <flatmm_sn_32x128x512_1x4x1_16x16x32.hpp>

Inheritance diagram for ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_BF16:
ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_Base

Public Types

using BDataType = bf16_t
 
using ODataType = bf16_t
 

Public Member Functions

template<typename BRes , typename BCoords , typename ORes , typename OCoords , typename OFlags , typename ScaleTensor >
CK_TILE_DEVICE auto operator() (const BRes &res_b, const BCoords &cached_coords_b, const ORes &res_o, const OCoords &cached_coords_o, const OFlags &o_flags, CK_TILE_LDS_ADDR void *smem, index_t n, const ScaleTensor &scale_, index_t tile_offset_b, index_t tile_offset_o)
 

Additional Inherited Members

- Static Public Member Functions inherited from ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_Base
static constexpr CK_TILE_DEVICE auto MakeCBlockDist ()
 
static constexpr CK_TILE_HOST_DEVICE ck_tile::index_t GetSmemSize ()
 
- Static Public Attributes inherited from ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_Base
static constexpr index_t Block_M = 32
 
static constexpr index_t Block_N = 128
 
static constexpr index_t Block_K = 512
 
static constexpr index_t WarpPerBlock_M = 1
 
static constexpr index_t WarpPerBlock_N = 4
 
static constexpr index_t WarpPerBlock_K = 1
 
static constexpr index_t Warp_M = 16
 
static constexpr index_t Warp_N = 16
 
static constexpr index_t Warp_K = 32
 
static constexpr index_t BlockSize = 256
 
static constexpr index_t Block_W = Warp_N * Warp_K
 
static constexpr index_t Block_Nr = Block_N / Warp_N
 
static constexpr index_t Block_Kr = Block_K / Warp_K
 
static constexpr index_t Repeat_M = Block_M / (Warp_M * WarpPerBlock_M)
 
static constexpr index_t Repeat_N = Block_N / (Warp_N * WarpPerBlock_N)
 
static constexpr index_t Repeat_K = Block_K / (Warp_K * WarpPerBlock_K)
 

Member Typedef Documentation

◆ BDataType

using ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_BF16::BDataType = bf16_t

◆ ODataType

using ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_BF16::ODataType = bf16_t

Member Function Documentation

◆ operator()()

template<typename BRes , typename BCoords , typename ORes , typename OCoords , typename OFlags , typename ScaleTensor >
CK_TILE_DEVICE auto ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_BF16::operator() ( const BRes &  res_b,
const BCoords &  cached_coords_b,
const ORes &  res_o,
const OCoords &  cached_coords_o,
const OFlags &  o_flags,
CK_TILE_LDS_ADDR void *  smem,
index_t  n,
const ScaleTensor &  scale_,
index_t  tile_offset_b,
index_t  tile_offset_o 
)
inline
The documentation for this struct was generated from the following file: