tilelang.rocm.intrinsics.mfma_layout
====================================

.. py:module:: tilelang.rocm.intrinsics.mfma_layout


Attributes
----------

.. autoapisummary::

   tilelang.rocm.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_m_n
   tilelang.rocm.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_n_k
   tilelang.rocm.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_n_m
   tilelang.rocm.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_k_n


Functions
---------

.. autoapisummary::

   tilelang.rocm.intrinsics.mfma_layout.shared_16x4_to_local_64x1_layout_A
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x1_to_16x4_layout_A
   tilelang.rocm.intrinsics.mfma_layout.shared_4x16_to_local_64x1_layout_B
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x1_to_4x16_layout_B
   tilelang.rocm.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_C
   tilelang.rocm.intrinsics.mfma_layout.shared_16x16_to_ldmatrix_64x4_layout
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x4_to_16x16_layout_A
   tilelang.rocm.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_A
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x4_to_16x16_layout_B
   tilelang.rocm.intrinsics.mfma_layout.shared_16x16_to_local_64x4_layout_B
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x4_to_16x16_layout_C_m_n
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x4_to_16x16_layout_C_n_m
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x8_to_16x32_layout_A
   tilelang.rocm.intrinsics.mfma_layout.shared_16x32_to_local_64x8_layout_A
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x8_to_16x32_layout_B
   tilelang.rocm.intrinsics.mfma_layout.shared_16x32_to_local_64x8_layout_B
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_16x64_layout_A
   tilelang.rocm.intrinsics.mfma_layout.shared_16x64_to_local_64x16_layout_A
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_16x64_layout_B
   tilelang.rocm.intrinsics.mfma_layout.shared_16x64_to_local_64x16_layout_B
   tilelang.rocm.intrinsics.mfma_layout.shared_32x32_to_local_64x16_layout_C
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_32x32_layout_C_n_m
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_32x32_layout_C_m_n
   tilelang.rocm.intrinsics.mfma_layout.shared_32x32_to_local_64x16_layout_A
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_32x32_layout_A
   tilelang.rocm.intrinsics.mfma_layout.shared_32x32_to_local_64x16_layout_B
   tilelang.rocm.intrinsics.mfma_layout.thread_id_shared_access_64x16_to_32x32_layout_B
   tilelang.rocm.intrinsics.mfma_layout.make_mfma_swizzle_layout


Module Contents
---------------

.. py:function:: shared_16x4_to_local_64x1_layout_A(i, j)

.. py:function:: thread_id_shared_access_64x1_to_16x4_layout_A(thread_id, local_id)

.. py:function:: shared_4x16_to_local_64x1_layout_B(i, j)

.. py:function:: thread_id_shared_access_64x1_to_4x16_layout_B(thread_id, local_id)

.. py:function:: shared_16x16_to_local_64x4_layout_C(i, j)

.. py:function:: shared_16x16_to_ldmatrix_64x4_layout(ind)

.. py:function:: thread_id_shared_access_64x4_to_16x16_layout_A(thread_id, local_id)

.. py:function:: shared_16x16_to_local_64x4_layout_A(i, j)

.. py:function:: thread_id_shared_access_64x4_to_16x16_layout_B(thread_id, local_id)

.. py:function:: shared_16x16_to_local_64x4_layout_B(i, j)

.. py:data:: shared_16x16_to_local_64x4_layout_m_n

.. py:data:: shared_16x16_to_local_64x4_layout_n_k

.. py:data:: shared_16x16_to_local_64x4_layout_n_m

.. py:data:: shared_16x16_to_local_64x4_layout_k_n

.. py:function:: thread_id_shared_access_64x4_to_16x16_layout_C_m_n(thread_id, local_id)

.. py:function:: thread_id_shared_access_64x4_to_16x16_layout_C_n_m(thread_id, local_id)

.. py:function:: thread_id_shared_access_64x8_to_16x32_layout_A(thread_id, local_id)

.. py:function:: shared_16x32_to_local_64x8_layout_A(i, j)

.. py:function:: thread_id_shared_access_64x8_to_16x32_layout_B(thread_id, local_id)

.. py:function:: shared_16x32_to_local_64x8_layout_B(i, j)

.. py:function:: thread_id_shared_access_64x16_to_16x64_layout_A(thread_id, local_id)

.. py:function:: shared_16x64_to_local_64x16_layout_A(i, j)

.. py:function:: thread_id_shared_access_64x16_to_16x64_layout_B(thread_id, local_id)

.. py:function:: shared_16x64_to_local_64x16_layout_B(i, j)

.. py:function:: shared_32x32_to_local_64x16_layout_C(i, j)

.. py:function:: thread_id_shared_access_64x16_to_32x32_layout_C_n_m(thread_id, local_id)

.. py:function:: 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``.


.. py:function:: shared_32x32_to_local_64x16_layout_A(i, j)

.. py:function:: thread_id_shared_access_64x16_to_32x32_layout_A(thread_id, local_id)

.. py:function:: shared_32x32_to_local_64x16_layout_B(i, j)

.. py:function:: thread_id_shared_access_64x16_to_32x32_layout_B(thread_id, local_id)

.. py:function:: make_mfma_swizzle_layout(shared_buf, vecSize=8)