For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /max/get-started.md).
Mojo struct
BlockScaledSmem
struct BlockScaledSmem[a_type: DType, b_type: DType, c_type: DType, sfa_dtype: DType, sfb_dtype: DType, transpose_b: Bool, *, config: BlockScaledMatmulConfig[a_type, b_type, c_type, sfa_dtype, sfb_dtype, transpose_b]]
SMEM struct for block-scaled matmul with CLC scheduler pipeline.
Thin wrapper over BlockScaledTileCore + SmemPipelineBundle.
Fields
- core (
BlockScaledSmem[a_type, b_type, c_type, sfa_dtype, sfb_dtype, transpose_b, config=config].Core): - pipelines (
BlockScaledSmem[a_type, b_type, c_type, sfa_dtype, sfb_dtype, transpose_b, config=config].Pipelines):
Implemented traits
comptime members
Core
comptime Core = BlockScaledTileCore[a_type, b_type, c_type, sfa_dtype, sfb_dtype, transpose_b, config=config]
Pipelines
comptime Pipelines = SmemPipelineBundle[(config // config), config.num_accum_pipeline_stages, config.num_clc_pipeline_stages, BlockScaledTilePayload[a_type, b_type, sfa_dtype, sfb_dtype, IndexList(config.block_tile_shape[Int(0)], config.block_tile_shape[Int(2)], __list_literal__=NoneType(None)), IndexList(config.block_tile_shape[Int(1)], config.block_tile_shape[Int(2)], __list_literal__=NoneType(None)), IndexList(sfa_dim0[config](), sfa_dim1[config](), __list_literal__=NoneType(None)), IndexList(sfb_dim0[config](), sfb_dim1[config](), __list_literal__=NoneType(None)), config.num_pipeline_stages]]
Methods
a_tiles
def a_tiles(ref[AddressSpace._value] self) -> Self.Core.ATileArray
Get A tile array accessor.
Returns:
Self.Core.ATileArray
b_tiles
def b_tiles(ref[AddressSpace._value] self) -> Self.Core.BTileArray
Get B tile array accessor.
Returns:
Self.Core.BTileArray
c_tiles
def c_tiles(ref[AddressSpace._value] self) -> Self.Core.CTileArray
Get C tile array accessor.
Returns:
Self.Core.CTileArray
sfa_tiles
def sfa_tiles(ref[AddressSpace._value] self) -> Self.Core.SFATileArray
Get SFA tile array accessor.
Returns:
Self.Core.SFATileArray
sfb_tiles
def sfb_tiles(ref[AddressSpace._value] self) -> Self.Core.SFBTileArray
Get SFB tile array accessor.
Returns:
Self.Core.SFBTileArray
ab_pipeline_size
static def ab_pipeline_size() -> Int
Total size of A+B tiles for all pipeline stages (in elements).
Returns:
sf_pipeline_size
static def sf_pipeline_size() -> Int
Total size of SFA+SFB tiles for all pipeline stages (in elements).
Returns:
c_output_size
total_tile_size
static def total_tile_size() -> Int
Total tile storage size (A+B+SFA+SFB+C) in elements.
Returns:
Was this page helpful?
Thank you! We'll create more content like this.
Thank you for helping us improve!