Return is control flow, but TIRx currently represents it as an
Evaluate-wrapped intrinsic call. This prevents return values from
participating naturally in statement traversal and requires special-case
handling across the pipeline.
This change introduces a reflected tirx.Return statement carrying an
Expr, wires it through TVMScript, statement visitors and mutators,
lowering, storage planning, and C/LLVM code generation, and removes the
legacy tirx.ret and T.ret surfaces.
This PR removes the suite-wide TIRx compute-capability gate and
localizes hardware skips to CUDA codegen and tile-primitive tests. It
keeps the original test parameterization unchanged, allowing parser,
printer, IR, transform, and other non-hardware TIRx tests to run in
regular CI while device-dependent cases are skipped when SM100 hardware
is unavailable. CUDA codegen helpers use explicit target architectures
where needed, and the run-only benchmark utility tests retain a local
SM100 gate.
This intentionally avoids adding separate compile/run parameter cases. A
follow-up PR can audit slow frontend execution tests and define a
focused runtime regression budget for real TIRx kernels.
Local validation:
- `pytest -n auto -m "not gpu" tests/python/tirx`: 486 passed, 79
skipped.
- `pytest -n auto -m gpu tests/python/tirx` without matching hardware:
1509 skipped.
- Pre-commit passed on all changed files.
## Rationale
`SizeVar` encodes nonnegativity in runtime subtype identity, which is
fragile under cloning and remapping. Symbolic integer values should use
one `Var` representation, with nonnegative facts recorded in the
analyzer at the use sites that establish them.
## Changes
- Remove `SizeVar` from the C++, Python, TE, TVMScript, FFI, visitor,
and serialization surfaces, and migrate callers to `Var`.
- Preserve the existing Relax constraint ownership model and use
`MarkGlobalNonNegValue` as the canonical path for global nonnegative
facts.
- Preserve `T.handle()` as the normal opaque-handle form. An optional
dtype constructs a typed pointer, with `T.handle("void")` reserved for
an explicit pointer-to-void.
## Summary
- Make `PrimExpr` a typed C++ view over `Expr` values whose
`ExprNode::ty` is `PrimType`, instead of using a separate runtime node
class as the proof of primitive-ness.
- Use the shared `ir::Call` node for Relax, TIRX, and primitive-valued
calls, while keeping primitive-only APIs explicit at their semantic
boundaries.
- Keep Python on the general `Expr` surface for primitive-typed values
so `isinstance` behavior does not imply a nominal primitive-expression
subclass.
## Design Rationale
The main advantage of this change is that common expression nodes such
as `Call` can be unified without specializing each one to `PrimType`. A
single `ir::Call` can represent a Relax tensor call, a Relax scalar
call, or a primitive-valued intrinsic call; the result type stored in
`ExprNode::ty` determines whether that particular value can be viewed as
`PrimExpr`.
This keeps the IR node hierarchy focused on expression structure rather
than result-type categories. Nodes that are intrinsically primitive,
such as integer and floating-point literals or TIRX primitive operators,
still have strongly typed C++ APIs and data structures. General nodes
whose result type may vary, such as `Call`, remain general `Expr` nodes
and are narrowed to `PrimExpr` only where primitive-only semantics are
required.
The PR also keeps the compatibility surface practical: C++
primitive-only APIs continue to accept `PrimExpr`, Python exposes a
compatibility predicate for checking the primitive typed category, and
visitors/printers use one natural `Call` path rather than duplicating
Relax and primitive call handling. Missing expression types are
represented explicitly with `Type::Missing()` so constructors can leave
type inference to later analysis without relying on nullable `Type`
values.
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
Validating the apache-tvm wheel in a minimal environment (no torch,
scipy, cloudpickle, or tornado installed) produced 33 pytest collection
errors from module-level imports of optional packages. Add
pytest.importorskip guards so these modules are reported as skipped
instead of erroring during collection.
Indirect import chains guarded:
- tvm.topi.testing imports scipy
- tvm.s_tir.meta_schedule.testing.local_rpc (tvm.rpc.tracker) requires
tornado
- tvm.s_tir.dlight.benchmark imports cloudpickle
Also remove a stray pre-license-header `import pytest` in
test_runtime_builtin_paged_attention_kv_cache_flashinfer.py.
## Summary
Follow-up work on top of the TIRx infrastructure bring-up (#19581). It
extends the TIRx operator-dispatch, codegen, and TVMScript surfaces with
the next batch of low-level programming features for Blackwell-class
GPUs, while keeping `s_tir` script support intact.
## Main Changes
- **op-dispatch**: warp `ldmatrix`/`stmatrix` copy dispatch; split CUDA
copy into register / gmem-smem / `ldgsts` paths; `tcgen05.ld/st`
`.16x{64,128,256}b` dispatch with a factory and M=128 layout;
element-wise broadcast at the layout level with a copy vec-alignment
fix.
- **gemm**: CUDA synchronous `mma.sync` tensor-core dispatch; accept a
Layout F C operand for M=64 MMAs.
- **op**: add the `permute_layout` primitive (replaces `permute_dims`).
- **tvmscript**: add the `Tx.jit` decorator, `Tx.constexpr` compile-time
params, and `Tx.wg_reg_tile`.
- **lower-tirx**: introduce the `Tx.device_entry()` marker (replacing
`ScopeKind::kKernel`); canonical thread filters that drop the
`Tx.filter` wrapper.
- **codegen**: add a typed-pointer byte-offset intrinsic; remove the
`entry_cluster_sync` codegen attribute.
## Validation
- `pre-commit run` (changed files) — clean
- `ninja -C build -j$(nproc)` — builds
- `pytest tests/python/tirx/ -n 16`
- `1997 passed, 39 skipped, 3 xpassed`
- `python -m pytest tests/python/all-platform-minimal-test`
- `37 passed, 105 skipped`
- `TVM_TEST_TARGETS=llvm pytest tests/python/tirx-analysis
tests/python/tirx-base tests/python/tirx-transform -n 16`
- `630 passed, 25 skipped, 8 xfailed, 1 xpassed`
## Local CI Notes
Several full CI-equivalent jobs are not locally reproducible because
this machine is missing parts of the Apache TVM CI environment (e.g.,
specific `llvm-config` versions, Vulkan, ROCm, ARM/QEMU cross-toolchain,
and web/wasm components). The Blackwell/Trainium kernel tests are
maintained downstream and are intentionally not part of this PR.
test_transform_lower_tirx.py imports and calls Simplify, but the pass is
named StmtSimplify (the only simplify pass exported from
tvm.tirx.transform). The stale name makes the module fail to import at
collection time. Use StmtSimplify so the test collects and runs.
## Background
The `tvm::ir` layer previously had a reverse dependency on
`tvm::script`, injected via the `TVM_OBJECT_ENABLE_SCRIPT_PRINTER()`
macro that added a `Script()` member method to IR node types (IRModule,
PrimExpr, Buffer, PrimFunc, Stmt). This violated the intended one-way
dependency: `script` should depend on `ir`, never the other way around.
Additionally, `PrinterConfigNode` accumulated dialect-specific fields
(`tir_prefix`, `tir_import_module`, `tirx_prefix`, `relax_prefix`) that
created leakage between the generic printer infrastructure and dialect
internals.
## Changes
This PR restores the clean dependency direction and encapsulates dialect
config properly, in 5 commits:
1. **Lift TVMScript entry point into `script/printer/printer.h`**: New
header `include/tvm/script/printer/printer.h` introduces:
- `tvm::Script()` free function replacing `TVMScriptPrinter::Script()`
static method
- `TVMScriptPrinter` class with vtable (`NodeFunctor<std::string(...)>`)
- `TVM_REGISTER_SCRIPT_AS_REPR` macro for registering per-type repr
callbacks
2. **Drop `TVM_OBJECT_ENABLE_SCRIPT_PRINTER` macro**: Remove the macro
from all IR headers (`ir/expr.h`, `ir/module.h`, `tirx/buffer.h`,
`tirx/function.h`, `tirx/stmt.h`), eliminating the reverse `ir` →
`script` dependency. All call sites of `.Script()` member methods
updated to use `tvm::Script()`.
3. **Move dialect-specific `PrinterConfig` fields to `extra_config`**:
Remove `tir_prefix`, `tir_import_module`, `tirx_prefix`, `relax_prefix`
from `PrinterConfigNode`. Dialect internals now read their config via
`GetExtraConfig<T>(key, fallback)` with dotted keys (e.g.,
`"tirx.prefix"`). `buffer_dtype` is kept as a top-level field alongside
`int_dtype`/`float_dtype` since it is a shared scalar-literal default,
not a dialect-specific knob.
4. **Python: drop dialect kwargs, expose `extra_config`**: Update
`PrinterConfig`, `Scriptable.script()`, `Scriptable.show()`,
`Scriptable._relax_script()`, and `BasePyModule.script()` to use
`extra_config: dict | None = None` instead of individual dialect kwargs.
The tirx auto-switch logic is preserved.
5. **Fix transitive include breakage**: Explicitly add direct includes
for `config.h` and `node_functor.h` where headers previously relied on
transitive paths through `expr.h`/`module.h`.
## Testing
- C++ unit tests: 118/118 pass
- TVMScript printer tests: 771 passed, 1 skipped, 1 xfailed
- TIR namespace tests
(`tests/python/tirx/test_printer_tir_namespaces.py`): 13/13 pass
- Relax AST printer tests: 24/24 pass
- Minimal platform tests: 37/37 pass
- Pre-commit (ASF headers, ruff, clang-format): all clean
## Summary
This PR adds the initial TIRx support needed for low-level programming
of Blackwell-class GPU architectures. As part of the ongoing TIRx
refactor, it introduces TVMScript support for directly scripting
advanced hardware features without relying on scheduling as the primary
programming interface.
The change keeps existing `s_tir` script support intact while making
direct scripting a first-class path for TIRx programs.
## Main Changes
- Add TIRx operator dispatch and layout infrastructure.
- Add TVMScript support for new low-level TIRx operations.
- Add analysis, transform, and lowering support for TIRx IR nodes.
- Add CUDA/Blackwell-oriented codegen and intrinsic coverage.
- Add Python and C++ integration points for TIRx scripting and runtime
support.
## Validation
- `pre-commit run --all-files`
- `ninja -C build -j32`
- `CUDA_VISIBLE_DEVICES=2 pytest tests/python/tirx/ -n 16`
- `1723 passed, 47 skipped, 32 warnings`
- `CUDA_VISIBLE_DEVICES=2 python -m pytest -v
tests/python/all-platform-minimal-test`
- `37 passed, 105 skipped`
- `TVM_TEST_TARGETS=llvm python -m pytest -v tests/python/tirx-analysis
tests/python/tirx-base tests/python/tirx-transform -n 16`
- `664 passed, 25 skipped, 9 xfailed, 1 xpassed`
## Local CI Notes
Some full CI-equivalent jobs were not locally reproducible because this
machine is missing parts of the Apache TVM CI environment, including
`llvm-config-15/17`, Vulkan, ROCm, Maven, Sphinx, Doxygen, Emscripten,
and ARM/QEMU cross-toolchain components. Metal-specific tests were
skipped locally because no Metal runtime is available.
Tianqi Chen