[Enhancement] Introduce pass_configs parameter for kernel Caching #452
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
* Added a new `pass_configs` parameter to the `tilelang.compile` function to allow for more flexible kernel compilation configurations. * Updated related classes and methods to accommodate the new parameter, ensuring compatibility across the codebase. * Enhanced the `torch_assert_close` function to include customizable tensor names for better debugging output. * Refactored input handling in example scripts to streamline the process of obtaining inputs for kernel execution.
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Jul 18, 2025
…le-ai#452) * [Enhancement] Introduce pass_configs parameter for kernel compilation * Added a new `pass_configs` parameter to the `tilelang.compile` function to allow for more flexible kernel compilation configurations. * Updated related classes and methods to accommodate the new parameter, ensuring compatibility across the codebase. * Enhanced the `torch_assert_close` function to include customizable tensor names for better debugging output. * Refactored input handling in example scripts to streamline the process of obtaining inputs for kernel execution. * lint fix
LeiWang1999
added a commit
to LeiWang1999/tilelang
that referenced
this pull request
Jul 20, 2025
…le-ai#452) * [Enhancement] Introduce pass_configs parameter for kernel compilation * Added a new `pass_configs` parameter to the `tilelang.compile` function to allow for more flexible kernel compilation configurations. * Updated related classes and methods to accommodate the new parameter, ensuring compatibility across the codebase. * Enhanced the `torch_assert_close` function to include customizable tensor names for better debugging output. * Refactored input handling in example scripts to streamline the process of obtaining inputs for kernel execution. * lint fix
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This pull request introduces several updates across multiple files to enhance kernel compilation, improve configurability, and refine tensor comparison utilities. The most significant changes include adding support for
pass_configsin kernel compilation, simplifying input generation in profiling, and improving tensor comparison error messages.Kernel Compilation Enhancements:
Added a
pass_configsparameter to several functions and methods, including_generate_key,cached, and various adapter-related methods, to allow more granular control over kernel compilation. This includes updates intilelang/cache/kernel_cache.py,tilelang/jit/adapter/ctypes/adapter.py,tilelang/jit/adapter/cython/adapter.py, andtilelang/jit/kernel.py. ([[1]](https://github.com/tile-ai/tilelang/pull/452/files#diff-9c43640a625dd7e0bc001b30e8d6d253d4d910ad3e83058ee5ff6014b6989d93R71),[[2]](https://github.com/tile-ai/tilelang/pull/452/files#diff-9c43640a625dd7e0bc001b30e8d6d253d4d910ad3e83058ee5ff6014b6989d93R97),[[3]](https://github.com/tile-ai/tilelang/pull/452/files#diff-9c43640a625dd7e0bc001b30e8d6d253d4d910ad3e83058ee5ff6014b6989d93L141-R145),[[4]](https://github.com/tile-ai/tilelang/pull/452/files#diff-2234f838318e2a6bc276679648d8f0baf10ededce1bf0a6443538fe96449ce74R122),[[5]](https://github.com/tile-ai/tilelang/pull/452/files#diff-fb6f070e77115039cffa90aea9d733c451c961c1c517ad15f5f58487a4ab1548R290),[[6]](https://github.com/tile-ai/tilelang/pull/452/files#diff-0eaba9f3529de6197aa7e13ba09aadf2253971a3571cf805678f4816e764e194R150),[[7]](https://github.com/tile-ai/tilelang/pull/452/files#diff-0eaba9f3529de6197aa7e13ba09aadf2253971a3571cf805678f4816e764e194R254),[[8]](https://github.com/tile-ai/tilelang/pull/452/files#diff-0eaba9f3529de6197aa7e13ba09aadf2253971a3571cf805678f4816e764e194R270),[[9]](https://github.com/tile-ai/tilelang/pull/452/files#diff-0eaba9f3529de6197aa7e13ba09aadf2253971a3571cf805678f4816e764e194R280))Modified the kernel compilation call in
examples/flash_decoding/example_mha_inference.pyto includepass_configsfor disabling TMA lowering, providing more flexibility in kernel behavior. ([examples/flash_decoding/example_mha_inference.pyL308-R308](https://github.com/tile-ai/tilelang/pull/452/files#diff-7263a27918434e2d64c23d90c8339d397aebc1090f2a6814742f4627f0c7b084L308-R308))Profiling and Input Simplification:
examples/linear_attention/example_retnet.pyby replacing manual tensor initialization with the profiler's_get_inputsmethod, streamlining the code. ([examples/linear_attention/example_retnet.pyL186-L211](https://github.com/tile-ai/tilelang/pull/452/files#diff-8f0f916b8916e34b1e6f7274b64e6ee123a2a4afe3dc90eca8fc4948a3056394L186-L211))Tensor Comparison Improvements:
Enhanced
torch_assert_closeintilelang/utils/tensor.pyto include customizable base and reference tensor names (base_nameandref_name) in error messages, improving clarity during debugging. ([[1]](https://github.com/tile-ai/tilelang/pull/452/files#diff-d53bf1f2e3725debd8d792d8d799bcd17ea37dd3f2019146d6e93b3a3f19be58R224-R225),[[2]](https://github.com/tile-ai/tilelang/pull/452/files#diff-d53bf1f2e3725debd8d792d8d799bcd17ea37dd3f2019146d6e93b3a3f19be58L299-R302))Updated
assert_allcloseintilelang/profiler/__init__.pyto passbase_nameandref_nameparameters totorch_assert_close, ensuring consistent naming conventions in tensor comparison. ([tilelang/profiler/__init__.pyR122-R123](https://github.com/tile-ai/tilelang/pull/452/files#diff-b3f17db8492573c451e8c620fb20ac43fee3e885f97135d7710ebd0512a9df70R122-R123))