tilelang.cuda.intrinsics.layout.mma_layout¶

Attributes¶

`shared_16x8_to_mma_32x4_layout_sr_a`
`shared_16x8_to_mma_32x4_layout_sr_b`
`shared_16x8_to_mma_32x4_layout_rs_a`
`shared_16x8_to_mma_32x4_layout_rs_b`
`shared_16x16_to_mma_32x8_layout_sr_a`
`shared_16x16_to_mma_32x8_layout_sr_b`
`shared_16x16_to_mma_32x8_layout_rs_a`
`shared_16x16_to_mma_32x8_layout_rs_b`
`shared_16x32_to_mma_32x16_layout_sr_a`
`shared_16x32_to_mma_32x16_layout_sr_b`
`shared_16x32_to_mma_32x16_layout_rs_a`
`shared_16x32_to_mma_32x16_layout_rs_b`
`shared_16x64_to_mma_32x32_layout_sr_a`
`shared_16x64_to_mma_32x32_layout_rs_a`
`shared_8x64_to_mma_32x16_layout_sr_b`
`shared_8x64_to_mma_32x16_layout_rs_b`

Functions¶

`ldmatrix_32x4_to_shared_16x8_layout_a`(thread_id, local_id)
`ldmatrix_32x4_to_shared_16x8_layout_b`(thread_id, local_id)
`ldmatrix_32x8_to_shared_16x16_layout`(thread_id, local_id)
`ldmatrix_trans_32x8_to_shared_16x16_layout`(thread_id, ...)
`ldmatrix_32x16_to_shared_16x32_layout_a`(thread_id, ...)
`ldmatrix_32x16_to_shared_16x32_layout_b`(thread_id, ...)
`metal_ct_store_32x16_to_16x32_layout`(thread_id, local_id)
`metal_ct_store_index_map`()
`mma_store_32x8_to_shared_16x16_layout`(thread_id, local_id)
`mma_store_32x2_to_shared_8x8_layout_fp64`(thread_id, ...)
`shared_16x8_to_mma_a_32x4_layout`(i, j)
`shared_16x8_to_mma_a_32x4_layout_trans`(i, j)
`shared_16x8_to_mma_b_32x4_layout`(i, j)
`shared_16x8_to_mma_b_32x4_layout_trans`(i, j)
`shared_16x16_to_mma_a_32x8_layout`(i, j)
`shared_16x16_to_mma_a_32x8_layout_trans`(i, j)
`shared_16x16_to_mma_b_32x8_layout`(i, j)
`shared_16x16_to_mma_b_32x8_layout_trans`(i, j)
`shared_16x32_to_mma_a_32x16_layout`(i, j)
`shared_32x16_to_mma_a_32x16_layout_trans`(i, j)
`shared_16x32_to_mma_b_32x16_layout`(i, j)
`shared_32x16_to_mma_b_32x16_layout_trans`(i, j)
`mma_32x8_to_shared_16x16_layout`(thread_id, local_id)
`mma_load_a_32x4_to_shared_16x8_layout`(thread_id, local_id)
`mma_load_b_32x4_to_shared_16x8_layout`(thread_id, local_id)
`mma_load_a_32x16_to_shared_16x32_layout`(thread_id, ...)
`mma_load_a_32x8_to_shared_16x16_layout`(thread_id, local_id)	groupID = %laneid >> 2
`mma_load_b_32x16_to_shared_16x32_layout`(thread_id, ...)
`mma_load_b_32x8_to_shared_16x16_layout`(thread_id, local_id)	groupID = %laneid >> 2
`shared_16x64_to_mma_a_32x32_layout`(i, j)	A fragment layout for m16n8k64 e2m1 row-major operand.
`shared_64x16_to_mma_a_32x32_layout_trans`(i, j)
`shared_8x64_to_mma_b_32x16_layout`(i, j)	B fragment layout for m16n8k64 e2m1 col-major operand.
`shared_64x8_to_mma_b_32x16_layout_trans`(i, j)
`ldmatrix_32x32_to_shared_16x64_layout_a`(thread_id)	Row-start addresses for ldmatrix.x4 loading A m16k64 e2m1.
`ldmatrix_32x32_to_shared_16x64_layout_b`(thread_id)	Row-start addresses for ldmatrix.x4 loading B n16k64 e2m1.
`ldmatrix_32x16_to_shared_8x64_layout_b`(thread_id)	Row-start addresses for ldmatrix.x2 loading B n8k64 e2m1.
`mma_load_a_32x32_to_shared_16x64_layout`(thread_id, ...)	Inverse: (thread_id, local_id) -> (m, k) for A m16k64 e2m1.
`mma_load_b_32x16_to_shared_8x64_layout`(thread_id, local_id)	Inverse: (thread_id, local_id) -> (n, k) for B n8k64 e2m1.
`shared_16x16_to_mma_32x8_smoothlayout`(i, j)
`shared_16x32_to_mma_32x16_smoothlayout`(i, j)
`shared_32x16_to_mma_32x16_smoothlayout`(i, j)
`get_swizzle_layout`(row_idx, col_idx, row_size, dtype)
`make_mma_swizzle_layout`(shared_buf[, is_smooth])

Module Contents¶

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_32x4_to_shared_16x8_layout_a(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_32x4_to_shared_16x8_layout_b(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_32x8_to_shared_16x16_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_trans_32x8_to_shared_16x16_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_32x16_to_shared_16x32_layout_a(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_32x16_to_shared_16x32_layout_b(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.metal_ct_store_32x16_to_16x32_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.metal_ct_store_index_map()¶

tilelang.cuda.intrinsics.layout.mma_layout.mma_store_32x8_to_shared_16x16_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.mma_store_32x2_to_shared_8x8_layout_fp64(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x8_to_mma_a_32x4_layout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x8_to_mma_a_32x4_layout_trans(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x8_to_mma_b_32x4_layout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x8_to_mma_b_32x4_layout_trans(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x8_to_mma_32x4_layout_sr_a¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x8_to_mma_32x4_layout_sr_b¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x8_to_mma_32x4_layout_rs_a¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x8_to_mma_32x4_layout_rs_b¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_a_32x8_layout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_a_32x8_layout_trans(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_b_32x8_layout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_b_32x8_layout_trans(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_32x8_layout_sr_a¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_32x8_layout_sr_b¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_32x8_layout_rs_a¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_32x8_layout_rs_b¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x32_to_mma_a_32x16_layout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_32x16_to_mma_a_32x16_layout_trans(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x32_to_mma_b_32x16_layout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_32x16_to_mma_b_32x16_layout_trans(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x32_to_mma_32x16_layout_sr_a¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x32_to_mma_32x16_layout_sr_b¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x32_to_mma_32x16_layout_rs_a¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x32_to_mma_32x16_layout_rs_b¶

tilelang.cuda.intrinsics.layout.mma_layout.mma_32x8_to_shared_16x16_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.mma_load_a_32x4_to_shared_16x8_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.mma_load_b_32x4_to_shared_16x8_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.mma_load_a_32x16_to_shared_16x32_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.mma_load_a_32x8_to_shared_16x16_layout(thread_id, local_id)¶

groupID = %laneid >> 2 threadID_in_group = %laneid % 4

row = groupID for ai where 0 <= i < 2 || 4 <= i < 6: groupID + 8 Otherwise

col = (threadID_in_group * 2) + (i & 0x1) for ai where i < 4 (threadID_in_group * 2) + (i & 0x1) + 8 for ai where i >= 4

tilelang.cuda.intrinsics.layout.mma_layout.mma_load_b_32x16_to_shared_16x32_layout(thread_id, local_id)¶

tilelang.cuda.intrinsics.layout.mma_layout.mma_load_b_32x8_to_shared_16x16_layout(thread_id, local_id)¶

groupID = %laneid >> 2 threadID_in_group = %laneid % 4

row = (threadID_in_group * 2) + (i & 0x1) for bi where i < 2: (threadID_in_group * 2) + (i & 0x1) + 8 for bi where i >= 2

col = groupID

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x64_to_mma_a_32x32_layout(i, j)¶: A fragment layout for m16n8k64 e2m1 row-major operand.

tilelang.cuda.intrinsics.layout.mma_layout.shared_64x16_to_mma_a_32x32_layout_trans(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_8x64_to_mma_b_32x16_layout(i, j)¶: B fragment layout for m16n8k64 e2m1 col-major operand.

tilelang.cuda.intrinsics.layout.mma_layout.shared_64x8_to_mma_b_32x16_layout_trans(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x64_to_mma_32x32_layout_sr_a¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x64_to_mma_32x32_layout_rs_a¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_8x64_to_mma_32x16_layout_sr_b¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_8x64_to_mma_32x16_layout_rs_b¶

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_32x32_to_shared_16x64_layout_a(thread_id)¶: Row-start addresses for ldmatrix.x4 loading A m16k64 e2m1.

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_32x32_to_shared_16x64_layout_b(thread_id)¶: Row-start addresses for ldmatrix.x4 loading B n16k64 e2m1.

tilelang.cuda.intrinsics.layout.mma_layout.ldmatrix_32x16_to_shared_8x64_layout_b(thread_id)¶: Row-start addresses for ldmatrix.x2 loading B n8k64 e2m1.

tilelang.cuda.intrinsics.layout.mma_layout.mma_load_a_32x32_to_shared_16x64_layout(thread_id, local_id)¶

Inverse: (thread_id, local_id) -> (m, k) for A m16k64 e2m1.

Matches CUTLASS/CuTe SM120 ALayout: Layout<Shape<Shape<_4,_8>, Shape<_8,_2,_2>>,

Stride<Stride<_128,_1>, Stride<_16,_8,_512>>>.

tilelang.cuda.intrinsics.layout.mma_layout.mma_load_b_32x16_to_shared_8x64_layout(thread_id, local_id)¶: Inverse: (thread_id, local_id) -> (n, k) for B n8k64 e2m1.

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x16_to_mma_32x8_smoothlayout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_16x32_to_mma_32x16_smoothlayout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.shared_32x16_to_mma_32x16_smoothlayout(i, j)¶

tilelang.cuda.intrinsics.layout.mma_layout.get_swizzle_layout(row_idx, col_idx, row_size, dtype, swizzle_bytes=None)¶

Parameters:: dtype (tvm.DataType | str)

tilelang.cuda.intrinsics.layout.mma_layout.make_mma_swizzle_layout(shared_buf, is_smooth=False)¶

Parameters:: is_smooth (bool)