[TIRx] Post-bringup follow-ups: op-dispatch, namespaces, launch bounds, gemm-async, backend reorg#19757
Conversation
There was a problem hiding this comment.
Code Review
This pull request restructures the TVM repository by moving target-specific backend codegen and runtime sources (such as CUDA, Hexagon, Metal, OpenCL, ROCm, Vulkan, and WebGPU) from src/target and src/runtime into a unified src/backend directory structure. It also reorganizes the corresponding Python bindings under python/tvm/backend and introduces a dynamic backend loading mechanism. The review feedback highlights a few critical issues: a missing analyzer variable instantiation in tcgen05.py that will cause a NameError at runtime, a bug in SMEMPool.commit where attr_frame.add_callback is incorrectly called on attr_frame itself instead of the parent frame (potentially causing a frame leak or deadlock), and an improvement opportunity in _AliasModule to override __setattr__ and __delattr__ to properly forward writes and deletes to the underlying aliased module.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
| @@ -516,80 +521,152 @@ def compute_canonical_params(buf, buf_region, dtype, is_transposed): | |||
| Tuple of (swizzle_mode, ldo, sdo, is_mn_major). | |||
| """ | |||
| region = list(buf_region.region) | |||
There was a problem hiding this comment.
The analyzer variable is used on line 614 but is not defined within the scope of compute_canonical_params, which will raise a NameError at runtime. Please instantiate analyzer = Analyzer() at the beginning of the function.
| region = list(buf_region.region) | |
| region = list(buf_region.region) | |
| analyzer = Analyzer() |
| if isinstance(attr_frame, frame_mod.AttrFrame): | ||
| from functools import partial | ||
|
|
||
| attr_frame.add_callback(partial(attr_frame.__exit__, None, None, None)) | ||
| attr_frame.__enter__() |
There was a problem hiding this comment.
In SMEMPool.commit, attr_frame.add_callback is called on attr_frame itself before it is entered. Since attr_frame is never exited manually, its callbacks are never run, leading to a frame leak/deadlock. It should be added to parent_frame instead.
| if isinstance(attr_frame, frame_mod.AttrFrame): | |
| from functools import partial | |
| attr_frame.add_callback(partial(attr_frame.__exit__, None, None, None)) | |
| attr_frame.__enter__() | |
| if isinstance(attr_frame, frame_mod.AttrFrame): | |
| from functools import partial | |
| parent_frame = frame_mod.current() | |
| parent_frame.add_callback(partial(attr_frame.__exit__, None, None, None)) | |
| attr_frame.__enter__() |
| def __getattr__(self, name: str): | ||
| return getattr(self.__dict__["__tvm_backend_module__"], name) | ||
|
|
||
| def __dir__(self): | ||
| return sorted(set(super().__dir__()) | set(dir(self.__dict__["__tvm_backend_module__"]))) |
There was a problem hiding this comment.
_AliasModule does not override __setattr__ and __delattr__, which causes attributes set on the aliased module to not be reflected in the underlying module. Please override them to forward writes and deletes to the underlying module.
def __getattr__(self, name: str):
return getattr(self.__dict__["__tvm_backend_module__"], name)
def __setattr__(self, name: str, value: Any):
if name in ("__tvm_backend_module__", "__package__", "__all__", "__path__"):
super().__setattr__(name, value)
else:
setattr(self.__dict__["__tvm_backend_module__"], name, value)
def __delattr__(self, name: str):
if name in ("__tvm_backend_module__", "__package__", "__all__", "__path__"):
super().__delattr__(name)
else:
delattr(self.__dict__["__tvm_backend_module__"], name)
def __dir__(self):
return sorted(set(super().__dir__()) | set(dir(self.__dict__["__tvm_backend_module__"])))"94e2474 to
d8aeb5b
Compare
…s, gemm-async, backend reorg Batch of post-bringup TIRx follow-ups, rebased onto current main: - Per-call exec scope via Tx.<scope>.op; remove ExecScopeStmt - Split TIRx op namespaces; remove tile-primitive kind attrs - Support explicit CUDA launch bounds - gemm-async: support contiguous-axis (K-major) operand slicing - Move in-tree GPU backends out of core into src/backend/<target>/ and python/tvm/backend/<target>/
d8aeb5b to
72a2111
Compare
This PR batches several post-bringup TIRx follow-ups, rebased onto current
main.Changes
Tx.<scope>.op; removeExecScopeStmtsrc/backend/<target>/andpython/tvm/backend/<target>/(codegen/runtime/op), with the correspondingCMakeLists.txt/cmake/modulesand include-path updatesTesting
USE_CUDA=ON/USE_LLVM=ONtests/python/tirx/) passes locally