* [feat] add lower-trace support for debugging
* [fix] address review: rebase, doc migration, lower_trace robustness
Address CodeRabbit + SiriusNEO review feedback on #2470:
Blocking (SiriusNEO):
- Rebase onto tile-ai/main (resolve docs/ + tilelang/__init__.py conflicts).
- Move IR Lower Trace doc from docs/tutorials/ to docs/tools/lower_trace.md
(alongside pass_diff.md); add index entry and 'superseded' pointer on
pass_diff.md; cross-link from the tutorial.
CodeRabbit Major:
- lower_trace/__init__.py: finally block uses active_exc pattern so a
generate_html() failure can no longer mask the real lowering exception.
- lower_trace/html.py: sanitize rec.phase into DOM-safe IDs via _safe_id()
before embedding in id="..." / selectors (avoids markup/JS breakage and
local HTML injection from special characters).
CodeRabbit Minor:
- test_lower_trace.py: restore caller's TL_LOWER_TRACE* env after the guarded
tilelang import instead of permanently clearing it.
- lower_trace/__init__.py: terminal mode prints a FAILED stanza (pass name,
exception, pre-pass IR) before re-raising so the failing step is visible.
- lower_trace/core.py: _get_codegen_output_path() no longer auto-writes
codegen.cpp in terminal-only mode (explicit override still honored).
- lower_trace/html.py: expandAll()/collapseCtx() null-guard the table lookup
so a changed diff with no table fails safe.
Cleanup:
- Add docstrings to test helpers and tools.__getattr__ (coverage now 94%+).
- ruff check + format clean.
* [Docs] Add lower trace to tools sidebar
* [fix] route lower-trace env vars through tilelang.env
Register TL_LOWER_TRACE / TL_LOWER_TRACE_DIR as EnvVar in the central
Environment class with get_lower_trace_mode()/get_lower_trace_dir()
parsers, mirroring the existing TILELANG_PASS_DIFF convention, instead
of reading os.environ directly in the lower_trace module.
* [fix] simplify lower_trace lazy-import in tools __getattr__
Drop the redundant globals() cache assignment (import_module already
binds the submodule onto the parent package) and fix the docstring: the
module is imported lazily for cost, not to defer monkey-patch side
effects, which happen in enable(). Note why importlib is required over
"from . import" (the latter re-enters __getattr__ and recurses).
---------
Co-authored-by: caojian5 <caojian5@huawei.com>
Co-authored-by: SiriusNEO <chaofan@deepseek.com>
* Add experimental IKET CUDA backend hooks
* Keep IKET APIs under namespace
* Use real Perfetto screenshot for IKET example
* Show IKET events in Perfetto timeline screenshot
* Add IKET validation overhead example
* Remove IKET validation overhead example
* Fix IKET lint issues
* Move IKET integration to CUDA tools
* Add tools documentation and fix workflow regressions
---------
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* [Feature] Add pass_visualizer: structure-tree pass browser
Add an interactive, pass-by-pass IR structure-tree visualizer under
tilelang/utils/pass_visualizer. It complements the existing text-level
pass_diff tool by rendering the SBlock structure tree and expanding tile
ops by field name, with per-class operator highlighting (tile op / sync /
lowered hardware intrinsic) across the CUDA lowering prologue.
Includes a CLI entry point, an example gemm_relu kernel, a pytest suite,
and a docs section contrasting it with pass_diff.
Co-Authored-By: Claude <noreply@anthropic.com>
* Move pass_visualizer to tilelang/tools and fix prologue + review issues
- Move package from tilelang/utils to tilelang/tools per review
- Add missing MaterializeKernelLaunch pass (fixes CUDA CI modulo-by-zero)
- Escape "</" in embedded JSON to prevent <script> injection in HTML report
- Use explicit utf-8 encoding for source/output file I/O
- Guard _fmt_shape against symbolic shapes and validate module spec
Co-Authored-By: Claude <noreply@anthropic.com>
---------
Co-authored-by: Claude <noreply@anthropic.com>
* [Feat] Add tilelang autodd for delta debugging
* fix typos
* fix lint error
* fix typos
* fix lint error
* fix bugs
* Apply suggestions from code review
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
* fix codeview comments
* [Refactor] Move AutoDD detection to env module and update import logic
* Refactor: Relocate the _is_running_autodd function to the env module for better organization and encapsulation.
* Update initialization logic to skip logger and heavy imports based on a new light import mode, enhancing flexibility in module usage.
* Ensure consistent handling of environment variables across the package, improving overall code clarity and maintainability.
* [Documentation] Add AutoDD section to debug_tools_for_tilelang.md
* Introduced a comprehensive guide on AutoDD (Automatic Delta Debugging) for isolating bugs in TileLang programs.
* Explained Delta Debugging methodology, usage, parameters, and provided examples for clarity.
* Highlighted the benefits of using AutoDD for large codebases and hard-to-locate errors, emphasizing time-saving aspects.
* Included tips for effective usage and a reference to a complete example in the documentation.
---------
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
Co-authored-by: kurisu6912 <227995639+kurisu6912@users.noreply.github.com>
Co-authored-by: LeiWang1999 <leiwang1999@outlook.com>
* Provide layout visualization tool
Adds a layout visualization tool to TileLang, which helps users understand and debug the layout transformations applied during compilation.
This tool visualizes the memory layout of tensors at different stages of the compilation process, allowing developers to identify potential inefficiencies and optimize their code for better performance.
The visualization can be enabled via a pass config option.
* format
* add layout visual example
* Adds vis extra with matplotlib dependency
* rafactor pass config name
* fix lint
* Enables configurable layout visualization formats
Allows users to specify the output formats (png, pdf, svg) for layout visualization through a pass config option.
This change provides more flexibility in how layout visualizations are generated, allowing users to choose the formats that best suit their needs.
It also fixes a bug where layout visualization was not correctly disabled when the config option was set to "false".
* Adds visual layout inference tool docs
* fix lint
* fix lint
* Rafactor configurable layout visualization formats
* fix lint
* fix typo
* add some comments
* fix lints
* add some warnings for user
* Moves layout visualization
* Refactors layout visualization pass configuration
Updates the layout visualization pass configuration to use boolean flag for enabling and a string for specifying formats.
* Enables multiple layout visualization formats
* Updates layout visualization docs
* Moves layout visualization to analysis
* [Refactor] Improve flash attention example and layout comparison logic
- Removed unnecessary annotation for `lse_local_split` in the flash attention example to streamline the code.
- Updated the handling of `lse_local_split` to utilize parallel processing for better performance.
- Refactored kernel compilation and profiling logic to enhance clarity and maintainability in the flash attention example.
- Added a condition in `FragmentNode::IsEqual` to handle broadcast cases, improving the robustness of layout comparisons.
* lint fix
* [Enhancement] Add support for shared memory scope in Fill operation
- Introduced handling for `shared.dyn` and `shared` memory scopes in the Fill operation.
- Implemented parallel operation and layout inference for improved performance in shared memory scenarios.
- Updated thread loop partitioning and vectorization logic to accommodate new memory scope handling.
* [Refactor] Remove deprecated decorator and enhance Cython kernel handling
- Removed the deprecated decorator from the main module and added a new implementation in the utils module for better organization.
- Introduced a pointer map in the Cython kernel adapter to manage pointer arguments, improving runtime shape resolution.
- Updated the Cython kernel wrapper to utilize the new pointer map for handling kernel arguments.
- Enhanced error checking in the tensor utility functions to ensure static shapes are enforced.
- Added a new proxy module for buffer and tensor handling, streamlining the interface for TIR programs.
* [Feature] Add matrix multiplication test and kernel implementation
- Introduced a new test file `test_tilelang_language_ptr.py` that implements a matrix multiplication function using TileLang's primitives.
- The `matmul_test` function defines a kernel for performing tile-level GEMM operations with customizable block sizes and data types.
- Added a `run_matmul` function to compile and execute the kernel, along with a test function to validate the implementation.
- Updated the `proxy.py` file to enhance type handling for buffer and tensor proxies, ensuring compatibility with TIR programs.
- Minor formatting improvements in `deprecated.py` for better readability.
* lint fix
* [Refactor] Update tensor creation in matrix multiplication test
- Replaced `T.Tensor.from_ptr` with `T.make_tensor` in `matmul_test` for improved clarity and consistency.
- Updated imports in `__init__.py` to include `make_tensor`.
- Added `make_tensor` function in `proxy.py` to streamline tensor creation from pointers.
* [Refactor] Update tensor definitions across multiple files
- Replaced instances of `T.Tensor` with updated tensor definitions in various benchmark and example files to enhance consistency and clarity.
- Adjusted tensor shapes and types in functions related to matrix multiplication, attention mechanisms, and other operations.
- Improved documentation in README and example files to reflect changes in tensor usage.
* lint fix
* [Refactor] Update tensor types in attention and matrix multiplication examples
- Replaced instances of `T.Tensor` with `T.SharedTensor` and `T.FragmentTensor` in various attention and matrix multiplication functions to improve consistency and clarity.
- Adjusted tensor definitions in benchmark and example files to align with the new tensor types.
- Enhanced the overall structure and readability of the code by standardizing tensor usage across multiple files.
* lint fix
* [Refactor] Update tensor types in GEMM example and test files
- Replaced instances of `T.Tensor` with `T.LocalTensor` and `T.Buffer` in the GEMM example and related test functions to improve consistency and clarity.
- Enhanced the overall structure of the code by standardizing tensor usage across multiple files, aligning with recent updates in tensor definitions.
* [Refactor] Update tensor usage in customize.py
- Replaced instances of `T.Tensor` with `T.Buffer` in the `reshape` and `view` functions to enhance consistency with recent tensor definitions.
- Improved code clarity by standardizing buffer usage across the file.
* [Refactor] Update tensor types in test_tilelang_transform_annotate_device_regions.py
- Replaced instances of `T.Tensor` with `T.Buffer` in the `before` and `expected` methods of the `TestAnnotateThreadExtent` and `TestAnnotateDeviceScope` classes to enhance consistency with recent tensor definitions.
- Improved code clarity by standardizing buffer usage across the test file.
* [Refactor] Update tensor types to SharedBuffer and FragmentBuffer
- Replaced instances of `T.SharedTensor` and `T.FragmentTensor` with `T.SharedBuffer` and `T.FragmentBuffer` across multiple benchmark, example, and test files to enhance consistency with recent tensor definitions.
- Improved code clarity and structure by standardizing buffer usage in attention and matrix multiplication functions.
* [Refactor] Introduce Tensor alias for Buffer in proxy.py
- Added a new alias `Tensor` for `Buffer` in `proxy.py` to facilitate JIT compilation, ensuring that inputs and outputs are mapped with `torch.Tensor`.
- This change enhances clarity and consistency in tensor usage across the codebase.
* Add GPU kernel for 2D continuous cumulative sum in TileLang example
- Introduced a new example script `example_tilelang_cumsum.py` that generates a GPU kernel for 2D continuous cumulative sum.
- Implemented functions to handle kernel configuration, memory allocation, and inclusive scan operations.
- Added a main execution block to demonstrate the kernel's functionality using PyTorch for tensor operations.
- Enhanced the example with error handling for power-of-two configurations and validation of results against PyTorch's built-in cumulative sum function.
* Refactor TileLang examples and enhance kernel compilation
- Updated `example_tilelang_cumsum.py` to improve GPU kernel generation for 2D continuous cumulative sum, including better parameter handling and error checking.
- Refactored `example_mha_bwd.py` to enhance kernel compilation readability and maintainability.
- Modified `kernel_cache.py` to prevent saving kernels to disk when using the DLPack backend, ensuring proper cache management.
- Added `get_block_bindings` function to `kernel.py` for improved access to block bindings in kernel launch frames.
- Cleaned up import statements in `__init__.py` for better organization and clarity.
* Enhance GPU kernel for 2D continuous cumulative sum in TileLang example
- Added additional spacing for improved readability in `example_tilelang_cumsum.py`.
- Refined kernel structure to enhance clarity and maintainability during GPU kernel generation for cumulative sum operations.
* Refactor CUDA post-processing callback registration in TileLang
- Introduced a new decorator `register_cuda_postproc_callback` for registering CUDA post-processing functions, enhancing usability and flexibility.
- Updated existing callback implementations to utilize the new decorator, improving code clarity and maintainability.
- Added debug prints to the CUDA code generation process for better traceability during development.
- Refactored the `OptimizeForTarget` function to streamline conditional statement handling in the pipeline transformation.
- Cleaned up the `inject_pipeline.cc` file by removing redundant code related to statement grouping and condition handling.
* lint fix
* Enhance BlockSparse GEMM Example with Autotuning and Configurable Parameters
- Added argument parsing to allow dynamic configuration of matrix dimensions and sparsity ratio.
- Implemented a function to generate various kernel configurations for autotuning.
- Refactored the main execution block to support both autotuned and default configurations.
- Improved the block mask generation to accommodate specified sparsity levels.
- Updated the kernel compilation process to utilize the new configurations and ensure accurate results verification.
pre-commit-ci[bot]