tilelang.intrinsics.wgmma_macro_generator
=========================================

.. py:module:: tilelang.intrinsics.wgmma_macro_generator


Attributes
----------

.. autoapisummary::

   tilelang.intrinsics.wgmma_macro_generator.lift


Classes
-------

.. autoapisummary::

   tilelang.intrinsics.wgmma_macro_generator.SwizzleMode
   tilelang.intrinsics.wgmma_macro_generator.TensorCoreIntrinEmitter


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

.. py:data:: lift

.. py:class:: SwizzleMode

   Bases: :py:obj:`enum.IntEnum`


   Enum where members are also (and must be) ints


   .. py:attribute:: NONE
      :value: 0



   .. py:attribute:: SWIZZLE_128B
      :value: 1



   .. py:attribute:: SWIZZLE_64B
      :value: 2



   .. py:attribute:: SWIZZLE_32B
      :value: 3



   .. py:method:: is_none()


   .. py:method:: is_swizzle_32b()


   .. py:method:: is_swizzle_64b()


   .. py:method:: is_swizzle_128b()


   .. py:method:: swizzle_byte_size()


   .. py:method:: swizzle_atom_size()


.. py:class:: TensorCoreIntrinEmitter(a_dtype = 'float16', b_dtype = 'float16', accum_dtype = 'float16', a_transposed = False, b_transposed = False, block_row_warps = 2, block_col_warps = 2, warp_row_tiles = 8, warp_col_tiles = 8, chunk = 16, reduce_k = 1, num_elems_per_byte = 1, is_m_first = False, thread_var = None)

   Bases: :py:obj:`tilelang.intrinsics.mma_macro_generator.TensorCoreIntrinEmitter`


   To eliminate Python syntax within TIR Macro.


   .. py:attribute:: wgmma_prefix
      :type:  str


   .. py:attribute:: a_shared_layout
      :type:  tilelang.layout.Layout
      :value: None



   .. py:attribute:: b_shared_layout
      :type:  tilelang.layout.Layout
      :value: None



   .. py:method:: wgmma(A_buf, B_buf, C_local_buf, clear_accum = False)


   .. py:method:: wgmma_rs(A_buf, B_buf, C_local_buf, clear_accum = False)


   .. py:method:: make_mma_load_layout(local_buf, matrix = 'A')

      Create a layout function for storing MMA results into a fragment buffer.
      This layout is used in conjunction with `inverse_mma_store_layout` to
      map fragment indices to threads and local indices.

      :param local_buf: The local buffer representing a fragment of a matrix.
      :type local_buf: tir.Buffer

      :returns: A fragment object that describes how threads and indices
                in `local_buf` are laid out.
      :rtype: T.Fragment

      :raises AssertionError: If `local_buf` is not detected to be a fragment buffer.



   .. py:method:: make_mma_store_layout(local_buf)

      Create a layout function for storing MMA results into a fragment buffer.
      This layout is used in conjunction with `inverse_mma_store_layout` to
      map fragment indices to threads and local indices.

      :param local_buf: The local buffer representing a fragment of a matrix.
      :type local_buf: tir.Buffer

      :returns: A fragment object that describes how threads and indices
                in `local_buf` are laid out.
      :rtype: T.Fragment

      :raises AssertionError: If `local_buf` is not detected to be a fragment buffer.