tilelang.cuda.pipeline
======================

.. py:module:: tilelang.cuda.pipeline


Attributes
----------

.. autoapisummary::

   tilelang.cuda.pipeline.cuda_pipeline


Functions
---------

.. autoapisummary::

   tilelang.cuda.pipeline.allow_warp_specialized
   tilelang.cuda.pipeline.module_has_tma
   tilelang.cuda.pipeline.CUDAPassPipelineBodyPrologue
   tilelang.cuda.pipeline.CUDAPassPipelineBody


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

.. py:function:: allow_warp_specialized(pass_ctx = None, target = None)

.. py:function:: module_has_tma(mod)

   Check if any function in the module was lowered with TMA operations.

   This reads the ``tl.has_tma`` attribute set by ``LowerTileOp`` during
   ``LowerAndLegalize``, which is the source of truth for whether TMA
   copies were actually generated.


.. py:function:: CUDAPassPipelineBodyPrologue(mod, target)

.. py:function:: CUDAPassPipelineBody(mod, target)

.. py:data:: cuda_pipeline