tilelang.ascend.language.print
==============================

.. py:module:: tilelang.ascend.language.print

.. autoapi-nested-parse::

   Ascend dialect T.print.

   Same debug externs as the common implementation (`debug_print_var`,
   `debug_print_buffer_value`, emitted by the Ascend codegen via
   src/tl_templates/ascend/debug.h), but with the NPU's execution model: a
   kernel body runs once per AI core and UB/L1 buffers are core-local, so
   buffer prints need no CUDA-style single-thread gating. Keeping the policy
   here (rather than probing hardware inside the CUDA dialect module) keeps
   each dialect self-contained.



Functions
---------

.. autoapisummary::

   tilelang.ascend.language.print.print


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

.. py:function:: print(obj = None, msg = '')

   Print a buffer, a primitive expression, or a plain message.

   :param obj: A ``tirx.Buffer`` (``local``, ``shared``, ``shared.dyn`` or
               ``global`` scope), a ``tirx.PrimExpr``, or ``None`` for a
               message-only print.
   :param msg: Optional message prefix.