FlatmmSn_32x128x512_1x4x1_16x16x32_BF16_itl Struct Reference

FlatmmSn_32x128x512_1x4x1_16x16x32_BF16_itl Struct Reference#

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

#include <flatmm_sn_32x128x512_1x4x1_16x16x32_itl.hpp>

Inheritance diagram for ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_BF16_itl:
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 CK_TILE_DEVICE constexpr auto MakeCBlockDist ()
static CK_TILE_HOST_DEVICE constexpr 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_itl::BDataType = bf16_t

◆ ODataType

using ck_tile::FlatmmSn_32x128x512_1x4x1_16x16x32_BF16_itl::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_itl::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: