tilelang.intrinsics.utils
=========================

.. py:module:: tilelang.intrinsics.utils


Functions
---------

.. autoapisummary::

   tilelang.intrinsics.utils.get_ldmatrix_offset
   tilelang.intrinsics.utils.shared_16x16_to_mma_32x8_layout
   tilelang.intrinsics.utils.shared_16x32_to_mma_32x16_layout
   tilelang.intrinsics.utils.shared_32x16_to_mma_32x16_layout
   tilelang.intrinsics.utils.mma_store_index_map
   tilelang.intrinsics.utils.mfma_store_index_map
   tilelang.intrinsics.utils.get_mma_micro_size
   tilelang.intrinsics.utils.index_to_coordinates


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

.. py:function:: get_ldmatrix_offset(matrix, row_idx, col_idx, stride, dtype = 'float16', transposed = False)

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

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

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

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

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

.. py:function:: get_mma_micro_size(dtype)

.. py:function:: index_to_coordinates(index, shape)

   General Implementation of:
       vjj = index % (micro_size_k // num_elems_per_byte)
       coordinates[-1] = index % shape[-1];
       vii = index // (micro_size_k // num_elems_per_byte) % micro_size_y
       index = index // shape[-1]; coordinates[-2] = index % shape[-2];
       vj = index // (micro_size_k // num_elems_per_byte * micro_size_y) % block_K // (micro_size_k // num_elems_per_byte)
       index = index // shape[-2]; coordinates[-3] = index % shape[-3];
       vi = index // (micro_size_k // num_elems_per_byte * micro_size_y * (block_K // (micro_size_k // num_elems_per_byte))) % block_N // micro_size_y
       index = index // shape[-3]; coordinates[-4] = index % shape[-4];