Skip to main content

/

DocsCode

Mojo struct

BlockwiseFP8TileCore

struct BlockwiseFP8TileCore[a_type: DType, b_type: DType, c_type: DType, a_scales_type: DType, transpose_b: Bool, *, config: MatmulConfig[a_type, b_type, c_type, transpose_b]]

Core tile storage for blockwise FP8 matmul SMEM structs.

Contains derived constants, layouts, tile storage, tile accessors, and size utilities. Shared between BlockwiseFP8Smem (CLC) and BlockwiseFP8_1D2DSmem (no CLC).

Fields​

  • ​tiles (BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].Tiles):

Implemented traits​

AnyType, ImplicitlyDestructible

comptime members​

AScalesTileArray​

comptime AScalesTileArray = BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].Tiles.AScalesTileArray

ATileArray​

comptime ATileArray = BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].Tiles.ATileArray

BK​

comptime BK = config.block_tile_shape[2]

BM​

comptime BM = config.block_tile_shape[0]

BN​

comptime BN = config.block_tile_shape[1]

BTileArray​

comptime BTileArray = BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].Tiles.BTileArray

CTileArray​

comptime CTileArray = BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].Tiles.CTileArray

MMA_M​

comptime MMA_M = config.mma_shape[0]

MMA_N​

comptime MMA_N = config.mma_shape[1]

num_accum_pipeline_stages​

comptime num_accum_pipeline_stages = config.num_accum_pipeline_stages

num_group_pipeline_stages​

comptime num_group_pipeline_stages = (BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].num_pipeline_stages // config)

num_output_stages​

comptime num_output_stages = config.num_output_stages

num_pipeline_stages​

comptime num_pipeline_stages = config.num_pipeline_stages

OutputM​

comptime OutputM = config.output_tile_shape[0]

OutputN​

comptime OutputN = config.output_tile_shape[1]

Payload​

comptime Payload = BlockwiseFP8TilePayload[a_type, b_type, a_scales_type, IndexList(BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BM, BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BK, __list_literal__=Tuple()), IndexList(BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BN, BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BK, __list_literal__=Tuple()), IndexList(1, BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BM, __list_literal__=Tuple()), BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].num_pipeline_stages]

Tiles​

comptime Tiles = BlockwiseFP8TileStorage[a_type, b_type, c_type, a_scales_type, IndexList(BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BM, BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BK, __list_literal__=Tuple()), IndexList(BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BN, BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BK, __list_literal__=Tuple()), BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].OutputM, BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].OutputN, IndexList(1, BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BM, __list_literal__=Tuple()), BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].num_pipeline_stages, BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].num_output_stages]

Methods​

a_tiles​

a_tiles(ref[AddressSpace._value] self) -> BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].ATileArray

Get A tile array accessor.

Returns:

BlockwiseFP8TileCore

b_tiles​

b_tiles(ref[AddressSpace._value] self) -> BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].BTileArray

Get B tile array accessor.

Returns:

BlockwiseFP8TileCore

c_tiles​

c_tiles(ref[AddressSpace._value] self) -> BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].CTileArray

Get C tile array accessor.

Returns:

BlockwiseFP8TileCore

a_scales_tiles​

a_scales_tiles(ref[AddressSpace._value] self) -> BlockwiseFP8TileCore[a_type, b_type, c_type, a_scales_type, transpose_b, config=config].AScalesTileArray

Get A-scales tile array accessor.

Returns:

BlockwiseFP8TileCore

ab_pipeline_size​

static ab_pipeline_size() -> Int

Total size of A+B tiles for all pipeline stages (in elements).

Returns:

Int

a_scales_pipeline_size​

static a_scales_pipeline_size() -> Int

Total size of A-scales tiles for all pipeline stages (in elements).

Returns:

Int

c_output_size​

static c_output_size() -> Int

Size of C tiles for all output stages (in elements).

Returns:

Int

total_tile_size​

static total_tile_size() -> Int

Total tile storage size (A+B+A-scales+C) in elements.

Returns:

Int

View source

View source

Was this page helpful?

Models

Cookbook

Puzzles

Blog

Community

Contact