tilelang.intrinsics.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

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, ...)
mma_store_32x8_to_shared_16x16_layout(thread_id, local_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_b_32x16_to_shared_16x32_layout(thread_id, ...)
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.intrinsics.mma_layout.ldmatrix_32x4_to_shared_16x8_layout_a(thread_id, local_id)

tilelang.intrinsics.mma_layout.ldmatrix_32x4_to_shared_16x8_layout_b(thread_id, local_id)

tilelang.intrinsics.mma_layout.ldmatrix_32x8_to_shared_16x16_layout(thread_id, local_id)

tilelang.intrinsics.mma_layout.ldmatrix_trans_32x8_to_shared_16x16_layout(thread_id, local_id)

tilelang.intrinsics.mma_layout.ldmatrix_32x16_to_shared_16x32_layout_a(thread_id, local_id)

tilelang.intrinsics.mma_layout.ldmatrix_32x16_to_shared_16x32_layout_b(thread_id, local_id)

tilelang.intrinsics.mma_layout.mma_store_32x8_to_shared_16x16_layout(thread_id, local_id)

tilelang.intrinsics.mma_layout.shared_16x8_to_mma_a_32x4_layout(i, j)

tilelang.intrinsics.mma_layout.shared_16x8_to_mma_a_32x4_layout_trans(i, j)

tilelang.intrinsics.mma_layout.shared_16x8_to_mma_b_32x4_layout(i, j)

tilelang.intrinsics.mma_layout.shared_16x8_to_mma_b_32x4_layout_trans(i, j)

tilelang.intrinsics.mma_layout.shared_16x8_to_mma_32x4_layout_sr_a

tilelang.intrinsics.mma_layout.shared_16x8_to_mma_32x4_layout_sr_b

tilelang.intrinsics.mma_layout.shared_16x8_to_mma_32x4_layout_rs_a

tilelang.intrinsics.mma_layout.shared_16x8_to_mma_32x4_layout_rs_b

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_a_32x8_layout(i, j)

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_a_32x8_layout_trans(i, j)

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_b_32x8_layout(i, j)

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_b_32x8_layout_trans(i, j)

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_32x8_layout_sr_a

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_32x8_layout_sr_b

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_32x8_layout_rs_a

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_32x8_layout_rs_b

tilelang.intrinsics.mma_layout.shared_16x32_to_mma_a_32x16_layout(i, j)

tilelang.intrinsics.mma_layout.shared_32x16_to_mma_a_32x16_layout_trans(i, j)

tilelang.intrinsics.mma_layout.shared_16x32_to_mma_b_32x16_layout(i, j)

tilelang.intrinsics.mma_layout.shared_32x16_to_mma_b_32x16_layout_trans(i, j)

tilelang.intrinsics.mma_layout.shared_16x32_to_mma_32x16_layout_sr_a

tilelang.intrinsics.mma_layout.shared_16x32_to_mma_32x16_layout_sr_b

tilelang.intrinsics.mma_layout.shared_16x32_to_mma_32x16_layout_rs_a

tilelang.intrinsics.mma_layout.shared_16x32_to_mma_32x16_layout_rs_b

tilelang.intrinsics.mma_layout.mma_32x8_to_shared_16x16_layout(thread_id, local_id)

tilelang.intrinsics.mma_layout.mma_load_a_32x4_to_shared_16x8_layout(thread_id, local_id)

tilelang.intrinsics.mma_layout.mma_load_b_32x4_to_shared_16x8_layout(thread_id, local_id)

tilelang.intrinsics.mma_layout.mma_load_a_32x16_to_shared_16x32_layout(thread_id, local_id)

tilelang.intrinsics.mma_layout.mma_load_b_32x16_to_shared_16x32_layout(thread_id, local_id)

tilelang.intrinsics.mma_layout.shared_16x16_to_mma_32x8_smoothlayout(i, j)

tilelang.intrinsics.mma_layout.shared_16x32_to_mma_32x16_smoothlayout(i, j)

tilelang.intrinsics.mma_layout.shared_32x16_to_mma_32x16_smoothlayout(i, j)

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

Parameters:

dtype (Union[tvm.DataType, str])

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

Parameters:

is_smooth (bool)