tilelang.cuda.language.kernel
=============================

.. py:module:: tilelang.cuda.language.kernel

.. autoapi-nested-parse::

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



Functions
---------

.. autoapisummary::

   tilelang.cuda.language.kernel.Kernel


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

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

   Construct a kernel launch frame for CUDA: a grid of thread blocks.

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

   :param \*blocks: Grid extent along each axis (1-3 dimensions, ``gridDim.(x|y|z)``). The
                    launch yields one block index per axis (``blockIdx.(x|y|z)``).
   :type \*blocks: int | PrimExpr
   :param threads: Threads per block: a count for ``blockDim.x`` or up to three
                   per-dimension extents for ``blockDim.(x|y|z)``. Defaults to 128 when
                   omitted.
   :type threads: int | list[int] | tuple[int, ...], optional
   :param prelude: CUDA source injected before the generated kernel, e.g. ``#include``
                   lines or helper functions.
   :type prelude: str, optional
   :param cluster_dims: Thread block cluster shape (SM90+). ``2`` or ``(2, 1, 1)`` launches
                        2-CTA clusters via ``cudaLaunchKernelEx``. ``T.ClusterKernel`` is the
                        same launch with a required ``cluster_dims``.
   :type cluster_dims: int | tuple[int, int, int] | list[int], optional

   .. rubric:: Examples

   .. code-block:: python

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

       with T.Kernel(grid_x, grid_y, threads=(64, 2)) as (bx, by):
           tx, ty = T.get_thread_bindings()
           ...