tilelang.intrinsics.mfma_layout

Attributes

shared_16x16_to_local_64x4_layout_m_n
shared_16x16_to_local_64x4_layout_n_k
shared_16x16_to_local_64x4_layout_n_m
shared_16x16_to_local_64x4_layout_k_n

Functions

shared_16x4_to_local_64x1_layout_A(i, j)
thread_id_shared_access_64x1_to_16x4_layout_A(...)
shared_4x16_to_local_64x1_layout_B(i, j)
thread_id_shared_access_64x1_to_4x16_layout_B(...)
shared_16x16_to_local_64x4_layout_C(i, j)
shared_16x16_to_ldmatrix_64x4_layout(ind)
thread_id_shared_access_64x4_to_16x16_layout_A(...)
shared_16x16_to_local_64x4_layout_A(i, j)
thread_id_shared_access_64x4_to_16x16_layout_B(...)
shared_16x16_to_local_64x4_layout_B(i, j)
thread_id_shared_access_64x4_to_16x16_layout_C_m_n(...)
thread_id_shared_access_64x4_to_16x16_layout_C_n_m(...)
thread_id_shared_access_64x8_to_16x32_layout_A(...)
shared_16x32_to_local_64x8_layout_A(i, j)
thread_id_shared_access_64x8_to_16x32_layout_B(...)
shared_16x32_to_local_64x8_layout_B(i, j)
thread_id_shared_access_64x16_to_16x64_layout_A(...)
shared_16x64_to_local_64x16_layout_A(i, j)
thread_id_shared_access_64x16_to_16x64_layout_B(...)
shared_16x64_to_local_64x16_layout_B(i, j)
shared_32x32_to_local_64x16_layout_C(i, j)
thread_id_shared_access_64x16_to_32x32_layout_C_n_m(...)
thread_id_shared_access_64x16_to_32x32_layout_C_m_n(...)	Return (m, n) = (row, col) for the 32x32 MFMA output register layout.
shared_32x32_to_local_64x16_layout_A(i, j)
thread_id_shared_access_64x16_to_32x32_layout_A(...)
shared_32x32_to_local_64x16_layout_B(i, j)
thread_id_shared_access_64x16_to_32x32_layout_B(...)
make_mfma_swizzle_layout(shared_buf[, vecSize])

Module Contents

tilelang.intrinsics.mfma_layout.shared_16x4_to_local_64x1_layout_A(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x1_to_16x4_layout_A(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_4x16_to_local_64x1_layout_B(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x1_to_4x16_layout_B(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_C(i, j)

tilelang.intrinsics.mfma_layout.shared_16x16_to_ldmatrix_64x4_layout(ind)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x4_to_16x16_layout_A(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_A(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x4_to_16x16_layout_B(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_B(i, j)

tilelang.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_m_n

tilelang.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_n_k

tilelang.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_n_m

tilelang.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_k_n

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x4_to_16x16_layout_C_m_n(thread_id, local_id)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x4_to_16x16_layout_C_n_m(thread_id, local_id)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x8_to_16x32_layout_A(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_16x32_to_local_64x8_layout_A(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x8_to_16x32_layout_B(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_16x32_to_local_64x8_layout_B(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_16x64_layout_A(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_16x64_to_local_64x16_layout_A(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_16x64_layout_B(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_16x64_to_local_64x16_layout_B(i, j)

tilelang.intrinsics.mfma_layout.shared_32x32_to_local_64x16_layout_C(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_32x32_layout_C_n_m(thread_id, local_id)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_32x32_layout_C_m_n(thread_id, local_id)

Return (m, n) = (row, col) for the 32x32 MFMA output register layout.

For v_mfma_i32_32x32x32_i8 (gfx950), each wave-64 lane holds 16 output i32 values. The column (N-dimension) is indexed by thread_id % 32 and the row (M-dimension) is given by the interleaved formula below. This function returns (m_idx, n_idx) matching the (row, col) convention expected by stmatrix.

tilelang.intrinsics.mfma_layout.shared_32x32_to_local_64x16_layout_A(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_32x32_layout_A(thread_id, local_id)

tilelang.intrinsics.mfma_layout.shared_32x32_to_local_64x16_layout_B(i, j)

tilelang.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_32x32_layout_B(thread_id, local_id)

tilelang.intrinsics.mfma_layout.make_mfma_swizzle_layout(shared_buf, vecSize=8)