tilelang.tools.lower_trace.core =============================== .. py:module:: tilelang.tools.lower_trace.core .. autoapi-nested-parse:: IR lower trace — zero-intrusion debug tool for visualizing compilation passes. Monkey-patches ``tvm.ir.transform.Pass.__call__`` and ``PassPipeline.lower`` to automatically capture IR before/after every pass and generate diff reports. This module reads its configuration (``TL_LOWER_TRACE`` / ``TL_LOWER_TRACE_DIR``) through the centralized ``tilelang.env`` environment, or via values passed programmatically to ``enable()``. Supports two architectures: - New: ``PassPipeline.lower`` (each backend registers a pipeline object) - Old: phase-based functions called from ``tilelang.engine.lower`` Usage:: TL_LOWER_TRACE=1 python my_kernel.py # HTML report TL_LOWER_TRACE=terminal python my_kernel.py # terminal diff only TL_LOWER_TRACE=both python my_kernel.py # both terminal and HTML Attributes ---------- .. autoapisummary:: tilelang.tools.lower_trace.core.STATUS_COMPLETED tilelang.tools.lower_trace.core.STATUS_FAILED tilelang.tools.lower_trace.core.STATUS_SKIPPED tilelang.tools.lower_trace.core.STATUS_CODEGEN Classes ------- .. autoapisummary:: tilelang.tools.lower_trace.core.LowerRecord Functions --------- .. autoapisummary:: tilelang.tools.lower_trace.core.enable tilelang.tools.lower_trace.core.disable tilelang.tools.lower_trace.core.reset Module Contents --------------- .. py:data:: STATUS_COMPLETED :value: 'completed' .. py:data:: STATUS_FAILED :value: 'failed' .. py:data:: STATUS_SKIPPED :value: 'skipped' .. py:data:: STATUS_CODEGEN :value: 'codegen' .. py:class:: LowerRecord Result of running a single pass. .. py:attribute:: phase :type: str .. py:attribute:: name :type: str .. py:attribute:: index :type: int .. py:attribute:: before_text :type: str .. py:attribute:: after_text :type: str .. py:attribute:: changed :type: bool .. py:attribute:: add_lines :type: int :value: 0 .. py:attribute:: del_lines :type: int :value: 0 .. py:attribute:: status :type: str :value: 'completed' .. py:attribute:: error_msg :type: str :value: '' .. py:function:: enable(*, mode=_UNSET, trace_dir=_UNSET, codegen_output=_UNSET) Enable IR pass tracing via monkey-patching. :param mode: Force a trace mode (``'terminal'``, ``'html'``, ``'both'``, or ``None`` to disable). When omitted, the mode is read from the ``TL_LOWER_TRACE`` env var (via ``tilelang.env``) or a prior ``enable`` override. :type mode: str | None, optional :param trace_dir: Force the trace output base directory. When omitted, falls back to the ``TL_LOWER_TRACE_DIR`` env var, then ``./tmp/lower_trace_dir``. :type trace_dir: str | None, optional :param codegen_output: Path to save the codegen-generated C++/CUDA/etc. source code. When omitted, defaults to ``/codegen.cpp`` (inside the per-script output directory, beside ``.run_records/``). Pass ``None`` explicitly to suppress all extra saves. See ``_wrap_codegen_ffi`` for the three-file (```` / ``.original`` / ``.latest``) patch-and-recompile workflow. :type codegen_output: str | None, optional .. py:function:: disable() Remove the pass tracing hook and restore original behavior. .. py:function:: reset() Clear collected records and section cache. ``_script_dir`` is preserved (stable across runs, holds codegen files + html symlink). ``_run_dir`` is also preserved: clearing it here would split a single run into two directories, because pre-pipeline passes (which lazily create it via ``_ensure_run_dir``) run before ``PassPipeline.lower``/``_wrap_phase`` invokes ``reset`` on its first run. A fresh ``_run_dir`` for each subsequent run is instead established directly in ``_traced_pipeline_lower`` / ``_wrap_phase`` (when ``_run_counter > 1``), without calling ``reset`` so that records keep accumulating across runs.