tilelang.layout.ascend¶

Ascend fractal/NZ layout wrappers (moved out of swizzle.py).

Classes¶

FractalLayoutInfo

Semantic view of a canonical Ascend fractal layout.

Functions¶

`make_ascend_nz_layout`(buffer)	Create an Ascend NZ (zN fractal) layout for a Cube operand buffer.
`make_ascend_compact_nz_layout`(buffer)	Create a compact staged-NZ layout for a pre-packed UB buffer.
`make_ascend_major_k_layout`(buffer[, k_align])	Create an Ascend fractal layout for logical `[.., MN, K]` tiles.
`make_ascend_major_mn_layout`(buffer[, k_align])	Create an Ascend fractal layout for logical `[.., K, MN]` tiles.
`make_ascend_l0c_layout`(buffer)	Create the fixed Ascend L0C accumulator fractal layout for logical
`make_ascend_sf_layout`(buffer)	Create an Ascend SF layout for MX scale-factor L1 buffers.
`make_strided_slice`(phys_buf, layout, logical_region)	Compute a 0-origin strided buffer view for a logical sub-region on a
`try_extract_fractal_layout`(layout, buffer)	Get semantic information from a canonical Ascend fractal layout.

Module Contents¶

tilelang.layout.ascend.make_ascend_nz_layout(buffer)¶

Create an Ascend NZ (zN fractal) layout for a Cube operand buffer.

A logical [.., rows, cols] tile maps to physical [.., rows / 16, cols / C0, 16, C0] – a grid of 16xC0 fractals – where C0 = 32 bytes / element_size is the Cube unit’s fractal width and the forward map is (.., r, c) -> [.., r / 16, c / C0, r % 16, c % C0].

Parameters:: buffer (tilelang._typing.BufferLikeType) – BufferLikeType (rank >= 2, constant last-2 dims)
Returns:: The NZ fractal layout for the buffer.
Return type:: Layout

tilelang.layout.ascend.make_ascend_compact_nz_layout(buffer)¶

Create a compact staged-NZ layout for a pre-packed UB buffer.

Unlike make_ascend_nz_layout(), this layout does not round the row extent of every leading stage up to 16. A logical [..., rows, cols] buffer is stored as consecutive [C0-group, rows, C0] blocks, with all leading dimensions flattened into the C0-group axis. The allocated row extent determines the pitch in 32-byte blocks, including any trailing gap rows. Partial copies retain that pitch.

This layout is intended for UB data that is produced directly in NZ order (for example with vsstb or a layout-driven T.copy) and then copied to L1. It is not an L1 Cube operand layout. For an ND-to-NZ T.copy, the allocation’s row extent includes any desired padding, for example (tile_rows + 1, cols). Copy only the valid region with T.copy(src, dst[:tile_rows, :]); gap rows are not initialized.

Parameters:: buffer (tilelang._typing.BufferLikeType)

tilelang.layout.ascend.make_ascend_major_k_layout(buffer, k_align=1)¶

Create an Ascend fractal layout for logical [.., MN, K] tiles.

The contiguous reduce-K axis is grouped by C0 and is the outer physical major: (.., mn, k) -> [.., k / C0, mn / 16, mn % 16, k % C0].

Parameters:

buffer (tilelang._typing.BufferLikeType)
k_align (int)

tilelang.layout.ascend.make_ascend_major_mn_layout(buffer, k_align=1)¶

Create an Ascend fractal layout for logical [.., K, MN] tiles.

The contiguous output M/N axis is grouped by C0 and is the outer physical major: (.., k, mn) -> [.., mn / C0, k / 16, k % 16, mn % C0].

Parameters:

buffer (tilelang._typing.BufferLikeType)
k_align (int)

tilelang.layout.ascend.make_ascend_l0c_layout(buffer)¶

Create the fixed Ascend L0C accumulator fractal layout for logical [.., M, N] tiles.

The physical shape is [.., N / 16, M / 16, 16, 16] and the forward map is (.., m, n) -> [.., n / 16, m / 16, m % 16, n % 16].

This layout can also annotate a UB buffer to preserve NZ in L0C-to-UB T.copy and T.dual_copy. Declare padding in the buffer shape; copy regions retain the physical panel pitch of that allocation and must cover whole 16-column panels. Buffer indices remain logical.

Parameters:: buffer (tilelang._typing.BufferLikeType)

tilelang.layout.ascend.make_ascend_sf_layout(buffer)¶

Create an Ascend SF layout for MX scale-factor L1 buffers.

A logical [..., MN, K] scale tile maps to physical [..., MN/16, K / (16 / dtype.bits), 16, 16 / dtype.bits]. For uint8/int8 SF this is [..., MN/16, K/2, 16, 2].

Parameters:: buffer (tilelang._typing.BufferLikeType)

tilelang.layout.ascend.make_strided_slice(phys_buf, layout, logical_region)¶

Compute a 0-origin strided buffer view for a logical sub-region on a layout-annotated physical buffer.

The result is a plain strided Buffer sharing phys_buf’s data pointer, with explicit shape (extents) and strides pointing at the sub-region’s base element, including phys_buf.elem_offset. No layout knowledge is needed by the consumer.

Parameters:

phys_buf – The physical (remapped) buffer, e.g. with NZ shape [M/16, K/C0, 16, C0].
layout – The Layout whose forward map describes the physical layout.
logical_region – One Range per layout input dim, [min, min+extent).

Returns:

A strided buffer view into phys_buf.

Return type:

Buffer

class tilelang.layout.ascend.FractalLayoutInfo(kind, c0, row_frac, rows, cols, outer0, outer1, c0_axis, row16_axis)¶

Semantic view of a canonical Ascend fractal layout.

This is a Python wrapper around the C++ TryExtractAscendFractalLayout utility. The layout kind is recorded by the canonical constructor; the remaining fields are derived from that tag and the buffer.

kind¶: 0 = MajorK, 1 = L0C (fixed 16×16), 2 = SF, 3 = MajorMN

c0¶: C0 value (or 16 for L0C, pack for SF)

row_frac¶: always 16

rows¶: logical rows (dim -2)

cols¶: logical cols (dim -1)

outer0¶: full physical extent of C0-axis outer for matrix data or row16-axis outer for SF

outer1¶: full physical extent of row16-axis outer for matrix data or C0-axis outer for SF

c0_axis¶: 0 = row, 1 = col (pack axis for SF)

row16_axis¶: 0 = row, 1 = col

__slots__ = ('kind', 'c0', 'row_frac', 'rows', 'cols', 'outer0', 'outer1', 'c0_axis', 'row16_axis')¶

kind¶

c0¶

row_frac¶

rows¶

cols¶

outer0¶

outer1¶

c0_axis¶

row16_axis¶

__repr__()¶

tilelang.layout.ascend.try_extract_fractal_layout(layout, buffer)¶

Get semantic information from a canonical Ascend fractal layout.

Returns:: FractalLayoutInfo, or None when the input is undefined/invalid.
Raises:: InternalError – If layout was not made by a canonical Ascend layout constructor or a transform dropped its semantic tag.