tilelang.layout.swizzle

Wrapping Layouts.

函式

make_swizzled_layout(buffer[, k_major, allow_pad])

make_volta_swizzled_layout(buffer[, is_a, k_inner])

make_wgmma_swizzled_layout(buffer[, continuity, k_major])

make_tcgen05mma_swizzled_layout(buffer[, continuity, ...])

make_full_bank_swizzled_layout(buffer)

make_half_bank_swizzled_layout(buffer)

make_quarter_bank_swizzled_layout(buffer)

make_linear_layout(buffer_or_load_or_region)

Create a row-major linear layout for any dimension.

make_gemm_fragment_8x8()

Create a standard 8x8 GEMM fragment layout for ldmatrix/stmatrix.

make_gemm_fragment_8x8_transposed()

Create a transposed 8x8 GEMM fragment layout for ldmatrix/stmatrix.

make_fully_replicated_layout_fragment(buffer, threads)

Create a fully replicated layout for a fragment buffer.

Module Contents

tilelang.layout.swizzle.make_swizzled_layout(buffer, k_major=True, allow_pad=True)
參數:

  • buffer (tilelang._typing.BufferLikeType)

  • k_major (bool)

  • allow_pad (bool)

tilelang.layout.swizzle.make_volta_swizzled_layout(buffer, is_a=True, k_inner=True)
參數:

  • buffer (tilelang._typing.BufferLikeType)

  • is_a (bool)

  • k_inner (bool)

tilelang.layout.swizzle.make_wgmma_swizzled_layout(buffer, continuity=None, k_major=True)
參數:

  • buffer (tilelang._typing.BufferLikeType)

  • continuity (int)

  • k_major (bool)

tilelang.layout.swizzle.make_tcgen05mma_swizzled_layout(buffer, continuity=None, k_major=True)
參數:

  • buffer (tilelang._typing.BufferLikeType)

  • continuity (int)

  • k_major (bool)

tilelang.layout.swizzle.make_full_bank_swizzled_layout(buffer)
參數:

buffer (tilelang._typing.BufferLikeType) -- BufferLikeType

範例

make_full_bank_swizzled_layout(buffer)

tilelang.layout.swizzle.make_half_bank_swizzled_layout(buffer)
參數:

buffer (tilelang._typing.BufferLikeType) -- BufferLikeType

範例

make_half_bank_swizzled_layout(buffer)

tilelang.layout.swizzle.make_quarter_bank_swizzled_layout(buffer)
參數:

buffer (tilelang._typing.BufferLikeType) -- BufferLikeType

範例

make_quarter_bank_swizzled_layout(buffer)

tilelang.layout.swizzle.make_linear_layout(buffer_or_load_or_region)

Create a row-major linear layout for any dimension.

參數:

buffer_or_load_or_region (tilelang._typing.BufferLikeType) -- BufferLikeType

回傳:

A row-major linear layout

回傳型別:

Layout

tilelang.layout.swizzle.make_gemm_fragment_8x8()

Create a standard 8x8 GEMM fragment layout for ldmatrix/stmatrix.

This layout matches the warp-level matrix multiplication pattern used in tensor cores.

回傳:

An 8x8 fragment layout

回傳型別:

Fragment

tilelang.layout.swizzle.make_gemm_fragment_8x8_transposed()

Create a transposed 8x8 GEMM fragment layout for ldmatrix/stmatrix.

This layout is the transposed version of make_gemm_fragment_8x8, useful for different access patterns in matrix operations.

回傳:

A transposed 8x8 fragment layout

回傳型別:

Fragment

tilelang.layout.swizzle.make_fully_replicated_layout_fragment(buffer, threads)

Create a fully replicated layout for a fragment buffer.

A fully replicated fragment means all threads hold identical copies of the entire buffer. This is useful for index buffers or masks that need to be accessed uniformly across all threads.

參數:

  • buffer (tilelang._typing.BufferLikeType) -- BufferLikeType to get shape information

  • threads (int) -- Number of threads (replicate extent)

回傳:

A fully replicated layout where each thread has a complete copy

回傳型別:

Fragment

範例

>>> C_local = T.alloc_fragment((2,), T.float32)
>>> layout = make_fully_replicated_layout_fragment(C_local, 256)
>>> T.annotate_layout({C_local: layout})