tilelang.metal.language.kernel
==============================

.. py:module:: tilelang.metal.language.kernel

.. autoapi-nested-parse::

   Metal dialect of ``T.Kernel``: the common launch plus Metal launch annotations.



Functions
---------

.. autoapisummary::

   tilelang.metal.language.kernel.Kernel


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

.. py:function:: Kernel(*blocks, threads = None, prelude = None)

   Construct a kernel launch frame for Metal: a grid of threadgroups.

   Code inside the launch operates at the threadgroup level: ``T.Parallel``,
   ``T.copy`` and friends are mapped onto threads by the compiler.
   ``T.get_thread_binding()`` exposes the thread index for thread-level code.
   The keyword arguments are recorded at trace time and materialized by the
   Metal pipeline once the target is known.

   :param \*blocks: Grid extent along each axis (1-3 dimensions). The launch yields one
                    threadgroup index per axis.
   :type \*blocks: int | PrimExpr
   :param threads: Threads per threadgroup: a count or up to three per-dimension extents.
                   Defaults to 128 when omitted.
   :type threads: int | list[int] | tuple[int, ...], optional
   :param prelude: Source injected before the generated kernel, e.g. ``#include`` lines or
                   helper functions.
   :type prelude: str, optional

   .. rubric:: Examples

   .. code-block:: python

       with T.Kernel(T.ceildiv(N, 128), threads=128) as bx:
           ...