Skip to content

[TIRx] Post-bringup follow-ups: op-dispatch, namespaces, launch bounds, gemm-async, backend reorg#19757

Merged
tqchen merged 1 commit into
apache:mainfrom
spectrometerHBH:apache-pr-tirx-followups-2
Jun 14, 2026
Merged

[TIRx] Post-bringup follow-ups: op-dispatch, namespaces, launch bounds, gemm-async, backend reorg#19757
tqchen merged 1 commit into
apache:mainfrom
spectrometerHBH:apache-pr-tirx-followups-2

Conversation

@spectrometerHBH

Copy link
Copy Markdown
Contributor

This PR batches several post-bringup TIRx follow-ups, rebased onto current main.

Changes

  • op-dispatch: per-call exec scope via Tx.<scope>.op; remove ExecScopeStmt
  • namespaces: split TIRx op namespaces; remove tile-primitive kind attrs
  • codegen: support explicit CUDA launch bounds
  • gemm-async: support contiguous-axis (K-major) operand slicing
  • backend reorg: move in-tree GPU backends out of core into src/backend/<target>/ and python/tvm/backend/<target>/ (codegen/runtime/op), with the corresponding CMakeLists.txt / cmake/modules and include-path updates

Testing

  • Builds with USE_CUDA=ON / USE_LLVM=ON
  • The TIRx Python test suite (tests/python/tirx/) passes locally

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

critical

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.

Suggested change
region = list(buf_region.region)
region = list(buf_region.region)
analyzer = Analyzer()

Comment on lines +525 to +529
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__()

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

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.

Suggested change
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__()

Comment on lines +56 to +60
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__"])))

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

high

_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__"])))"

@spectrometerHBH spectrometerHBH force-pushed the apache-pr-tirx-followups-2 branch 5 times, most recently from 94e2474 to d8aeb5b Compare June 13, 2026 23:10
…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>/
@spectrometerHBH spectrometerHBH force-pushed the apache-pr-tirx-followups-2 branch from d8aeb5b to 72a2111 Compare June 13, 2026 23:50
@tqchen tqchen merged commit bb6f8ae into apache:main Jun 14, 2026
8 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants