Flatmm_32x512x128_1x4x1_16x16x32_Base Struct Reference

Flatmm_32x512x128_1x4x1_16x16x32_Base Struct Reference#

Composable Kernel: ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base Struct Reference
ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base Struct Reference

#include <flatmm_32x512x128_1x4x1_16x16x32.hpp>

Inheritance diagram for ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base:
ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_BF16 ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_FP16

Static Public Member Functions

static CK_TILE_DEVICE constexpr auto MakeCBlockDist ()
static CK_TILE_DEVICE constexpr auto MakeCBlockTile ()
static CK_TILE_HOST_DEVICE constexpr auto MakeLdsStoreDesc_A ()
static CK_TILE_HOST_DEVICE constexpr auto MakeLdsLoadDesc_A ()
static constexpr auto GetGemm_AWarpEnc ()
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSize ()

Static Public Attributes

static constexpr index_t Block_M = 32
static constexpr index_t Block_N = 512
static constexpr index_t Block_K = 128
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 NumWarps = 4
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 SubKPacks = 2
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 Function Documentation

◆ GetGemm_AWarpEnc()

constexpr auto ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::GetGemm_AWarpEnc ( )
inlinestaticconstexpr

◆ GetSmemSize()

CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::GetSmemSize ( )
inlinestaticconstexpr

◆ MakeCBlockDist()

CK_TILE_DEVICE constexpr auto ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::MakeCBlockDist ( )
inlinestaticconstexpr

◆ MakeCBlockTile()

CK_TILE_DEVICE constexpr auto ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::MakeCBlockTile ( )
inlinestaticconstexpr

◆ MakeLdsLoadDesc_A()

CK_TILE_HOST_DEVICE constexpr auto ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::MakeLdsLoadDesc_A ( )
inlinestaticconstexpr

◆ MakeLdsStoreDesc_A()

CK_TILE_HOST_DEVICE constexpr auto ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::MakeLdsStoreDesc_A ( )
inlinestaticconstexpr

Member Data Documentation

◆ Block_K

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Block_K = 128
staticconstexpr

◆ Block_Kr

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Block_Kr = Block_K / Warp_K
staticconstexpr

◆ Block_M

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Block_M = 32
staticconstexpr

◆ Block_N

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Block_N = 512
staticconstexpr

◆ Block_Nr

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Block_Nr = Block_N / Warp_N
staticconstexpr

◆ Block_W

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Block_W = Warp_N * Warp_K
staticconstexpr

◆ BlockSize

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::BlockSize = 256
staticconstexpr

◆ NumWarps

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::NumWarps = 4
staticconstexpr

◆ Repeat_K

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Repeat_K = Block_K / (Warp_K * WarpPerBlock_K)
staticconstexpr

◆ Repeat_M

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Repeat_M = Block_M / (Warp_M * WarpPerBlock_M)
staticconstexpr

◆ Repeat_N

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Repeat_N = Block_N / (Warp_N * WarpPerBlock_N)
staticconstexpr

◆ SubKPacks

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::SubKPacks = 2
staticconstexpr

◆ Warp_K

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Warp_K = 32
staticconstexpr

◆ Warp_M

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Warp_M = 16
staticconstexpr

◆ Warp_N

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::Warp_N = 16
staticconstexpr

◆ WarpPerBlock_K

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::WarpPerBlock_K = 1
staticconstexpr

◆ WarpPerBlock_M

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::WarpPerBlock_M = 1
staticconstexpr

◆ WarpPerBlock_N

index_t ck_tile::Flatmm_32x512x128_1x4x1_16x16x32_Base::WarpPerBlock_N = 4
staticconstexpr

The documentation for this struct was generated from the following file: