tvm.testing's test-gating layer had a number of one-line helper
functions that add a name but no behavior. Inline the thin ones so call
sites name the underlying flag/feature/condition directly.
Pytest plugin (plugin.py): _target_to_requirement built its skip / gpu
marks through two one-line wrappers (_gpu_mark_and_skip / _skip_only)
plus a per-kind if ladder. Replace them with two frozensets (GPU- vs
CPU-family kinds) and resolve the skip probe by name:
marks.append(pytest.mark.skipif(not getattr(env, f"has_{kind}")(),
reason=f"need {kind}"))
The cuda+cudnn / cuda+cublas accelerator-library cases are remapped
inline (cudnn before cublas). Adds two direct unit tests for the
cudnn/cublas special-case and the unknown-kind ([]) fallback.
tvm.testing.env (env.py): inline the pure probe wrappers that just
forwarded to a primitive --
* build-flag (5): has_cutlass/rpc/nnapi/openclml/mrvl ->
env.build_flag_enabled ("USE_X"). The private _build_flag_enabled is
promoted to the public build_flag_enabled; the composed probes
(has_cudnn/cublas/nccl/hipblas) and the hexagon/adreno probes call it
too.
* cpu-feature (5 pure):
has_arm_dot/arm_fp16/aarch64_sve/aarch64_sme/x86_amx ->
env.has_cpu_feature("..."). The composed has_x86_vnni (avx512vnni OR
avxvnni) and has_x86_avx512 (a five-feature set) are kept -- not thin
wrappers.
Also drops the obsolete test_build_flag_probe_matches_libinfo self-test
and the matching _BOOL_PROBES entries.
The runtime device probes (has_cuda/has_rocm/...) are intentionally left
as-is: the pytest plugin resolves env.has_<kind>() from each target
kind, so those names are load-bearing rather than thin wrappers.
This pr modernizes test gating. It replaces the heavy
`tvm.testing.Feature` machinery with a thin `tvm.testing.env` module of
`has_*()` capability probes, used via standard pytest.mark + skipif. And
markers move to `pyproject.toml`
This PR adds an autoload mechanism for out-of-tree backends, simplifies
TVM's Python library loading, and removes `TVMError` in favor of native
Python errors.
## Autoload out-of-tree backends
Out-of-tree packages can register an autoload callable under the
`tvm.backends` entry-point group (mirroring torch's device-backend
autoload). At `import tvm` startup each entry point is discovered and
its callable invoked once, after the core runtime and the `tvm`
namespace are fully initialized, so an extension can register
ops/targets/funcs or load extra libraries.
```toml
[project.entry-points."tvm.backends"]
tvm_foo = "tvm_foo:_autoload"
```
A failing extension is caught and surfaced via `warnings.warn` so it
cannot break `import tvm`. Autoload can be disabled with
`TVM_DEVICE_BACKEND_AUTOLOAD=0`.
## Simplify library loading
The library-loading path in `base.py` is consolidated around a single
`_LOADED_LIBS` dict (basename to ctypes handle) so downstream and
autoloaded extensions can skip already-loaded libraries; the per-backend
runtime DSO list is folded into `load_backend_libs`. Accumulated cruft
is removed: the Python-3.9 check, the readline shim, the `_FFI_MODE`
ctypes check, the `base.__version__` re-export, and `py_str` (call sites
inline `.decode("utf-8")`).
## Remove TVMError in favor of native Python errors
`TVMError` added a layer atop `RuntimeError` that downstream code had to
import and learn. It is removed; the registered FFI error kinds
(`InternalError`, `RPCError`, `OpError`, `DiagnosticError`,
`ScheduleError`) now subclass `RuntimeError` directly while staying
registered, so the FFI keeps throwing the right kinds. All `TVMError`
imports, `except`/`raise`/`isinstance` uses, and
`pytest.raises(tvm.TVMError)` sites move to the `RuntimeError` builtin.
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
Lifts 10 host-toolchain / CLI / process / utility modules from
`python/tvm/contrib/` to a new `python/tvm/support/` package, and
deletes two dead contrib shims.
`tvm.support` is the home for Python helpers that integrate TVM with
external CLIs and host-side tools — compilers, archivers, subprocess
pools, and build-info queries. These are load-bearing internal pieces
that TVM's compile/link/run paths depend on. `tvm.contrib` is reserved
for optional vendor SDK integrations and experimental features. The
distinction is documented in the `tvm.support` package docstring.
Moved (one commit each):
- `tvm.contrib.cc` → `tvm.support.cc`
- `tvm.contrib.nvcc` → `tvm.support.nvcc`
- `tvm.contrib.rocm` → `tvm.support.rocm`
- `tvm.contrib.ndk` → `tvm.support.ndk`
- `tvm.contrib.xcode` → `tvm.support.xcode`
- `tvm.contrib.clang` → `tvm.support.clang`
- `tvm.contrib.emcc` → `tvm.support.emcc`
- `tvm.contrib.popen_pool` → `tvm.support.popen_pool`
- `tvm.contrib.utils` → `tvm.support.utils`
- `tvm.contrib.tar` → `tvm.support.tar`
Deleted:
- `tvm.contrib.spirv` — single `optimize()` wrapping `spirv-opt`; zero
importers.
- `tvm.contrib.rpc` — self-deprecation shim with "removed in 0.5"
banner; honoring it.
Package conversion:
- `python/tvm/support.py` → `python/tvm/support/__init__.py` with
inclusion-rule docstring.
- `libinfo()` extracted into `python/tvm/support/libinfo.py`.
- `FrontendTestModule` dropped (audit confirmed zero callers outside its
own definition).
## Compatibility
Hard break — no `tvm.contrib.<mod>` re-export shims. All callers updated
in this PR.
C++-side FFI registry keys (`tvm.contrib.nvcc.*`, etc.) are unchanged —
only the Python module path moves. Renaming the FFI keys is a separate
follow-up.
## 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.
## Motivation
Historically TVM ships a single monolithic `libtvm.so` that bundles both
the
runtime and the compiler/LLVM-heavy code paths. Deployment scenarios
that only
need the runtime end up paying the full compiler footprint (LLVM-static
dominates
the binary size), and the layout makes it awkward to install the project
under a
single Python package directory the way
`tvm_ffi`/`libinfo.load_lib_ctypes`
expects.
This PR splits the single shared library into two:
- `libtvm_runtime.so` — runtime-only symbols (loaded `RTLD_GLOBAL`).
- `libtvm_compiler.so` — compiler / LLVM / codegen, links
`libtvm_runtime.so`
publicly (loaded `RTLD_LOCAL`).
## Target restructure
- New CMake target `tvm_compiler` replaces the old `tvm` SHARED target.
- `tvm_compiler` depends on `tvm_runtime` via `target_link_libraries(...
PUBLIC tvm_runtime)`,
so anything that linked the old `tvm` now picks up the runtime
transitively.
- `tvm_libinfo_objs` (build-info TU) moved from `tvm_runtime` into
`tvm_compiler`
— it is compiler-side metadata and the runtime no longer needs it.
- All `target_link_libraries` / `target_compile_*` /
`set_target_properties` /
`tvm_ffi_add_apple_dsymutil` callsites have been rewired.
- The separate `libtvm_allvisible.so` target is **removed** (was only
consumed
by cpptests). Cpptests with private-symbol deps are deleted; remaining
cpptests now link directly against `libtvm_compiler.so` /
`libtvm_runtime.so`. `src/support/hexdump.cc` is folded into the header.
- `BUILD_DUMMY_LIBTVM` and the `BUILD_FOR_HEXAGON + USE_HEXAGON_GTEST`
cpp-test wiring are removed.
## Output and install layout
- All artifacts now go to `build/lib/` (was `build/`):
- `build/lib/libtvm_runtime.so`
- `build/lib/libtvm_compiler.so`
- Install layout is now `<package>/lib/` so
`tvm_ffi.libinfo.load_lib_ctypes`
with `package="tvm"` finds the libs in the wheel.
- CI Jenkins stash paths and `apps/hexagon_*` paths updated to the new
`build/lib/...` location.
## Python loader change
`python/tvm/base.py` now resolves the libs directly via a small
`package_lib_paths()` helper in `python/tvm/libinfo.py` (anchored on
`python/tvm/__file__`, returning the wheel `lib/`,
`<worktree>/build/lib`, and
`<worktree>/lib` candidates). Module-level `_LIB_RUNTIME`, `_LIB`, and
`_RUNTIME_ONLY` are set inline at import time:
- `libtvm_runtime.{so,dylib,dll}` loaded `RTLD_GLOBAL`.
- `libtvm_compiler.{so,dylib,dll}` loaded `RTLD_LOCAL`.
- `TVM_USE_RUNTIME_LIB` (parsed strictly: `1`/`true`/`yes`) selects
runtime-only at the loader level.
- When the compiler lib is absent, `_RUNTIME_ONLY` is set to True
automatically and `_LIB is _LIB_RUNTIME`.
## Non-obvious build-integration fixes
Three issues surfaced once both libs are loaded into the same process
and are
worth calling out:
1. **`fpA_intB_gemm` double-registration.** `fpA_intB_gemm_tvm` is an
OBJECT
library that registers a global `fastertransformer.gemm_fp16_int` at
static
init. Linking it into both `tvm_runtime` and `tvm_compiler` made the
registration run twice and trip the duplicate-registration check. Fix:
link
it (and the other runtime-only externals — `flash_attn`, NCCL, NVSHMEM,
RCCL) only into `tvm_runtime`. `tvm_compiler` picks them up via the
PUBLIC
`tvm_runtime` link.
2. **`-Wl,--no-as-needed` for minrpc.** `python/tvm/rpc/minrpc.py`
defaults
to `runtime="libtvm_runtime"` and passes `-Wl,--no-as-needed` so the
runtime static initializers actually run in the spawned minrpc binary
(without it, the linker drops the lib because no symbol is referenced
directly from the minrpc TU). minrpc does **not** link
`libtvm_compiler.so`.
3. **`testing.GetShape{Elem,Size}` moved to runtime.** Those two test
helpers
(the only `testing.*` symbols the minrpc test exercises) were registered
in
`src/support/ffi_testing.cc` (compiler-side). They are now registered in
`src/runtime/rpc/testing.cc` under `rpc.testing.GetShape{Elem,Size}` so
the minrpc server binary — runtime-only — can resolve them.
## Deprecations and breaking changes
- `BUILD_DUMMY_LIBTVM` is **removed** (option, libinfo entry, and CMake
wiring). Downstream consumers that built the dummy variant should link
`libtvm_runtime.so` directly.
- **Breaking change for downstream consumers** that read `libtvm.so` by
name:
there is no longer a `libtvm.so`. Replace with `libtvm_compiler.so`
(full)
or `libtvm_runtime.so` (runtime-only). The Vulkan device comment and a
few
test/CI comments have been updated accordingly.
- `libtvm_allvisible.so` is **removed**. Cpptests that depended on
private
out-of-line symbols have been deleted; the remaining cpp-test contract
is
documented as "public API or private header-only API only" (see
`tests/cpp/`).
- `tests/cpp-runtime/` (Hexagon + OpenCL backend tests) is **removed**
until
TVM moves to a plugin-mode backend architecture where each backend can
ship its own test harness with its own visibility scope.
## Tested
- `ninja` build: `build/lib/libtvm_runtime.so`,
`build/lib/libtvm_compiler.so`;
no `build/libtvm.so`, no `build/lib/libtvm_allvisible.so`.
`ldd build/lib/libtvm_compiler.so` links `libtvm_runtime.so`,
`libtvm_ffi.so`, `libfpA_intB_gemm.so`, `libflash_attn.so`.
- `ldd build/cpptest`: only `libtvm_compiler.so` + `libtvm_runtime.so` +
`libtvm_ffi.so` (no `libtvm_allvisible.so`).
- `./build/cpptest`: 144 / 144 tests pass across 29 suites.
- Smoke imports: full and `TVM_USE_RUNTIME_LIB=1` — both pass.
`TVM_USE_RUNTIME_LIB=0` correctly disables runtime-only mode (strict
parse).
- `tests/python/all-platform-minimal-test`: 75 passed, 77 skipped.
- `tests/python/runtime/`: 81 passed, 2 skipped (incl.
`test_rpc_return_remote_object` exercising the minrpc executable
end-to-end
via `rpc.testing.GetShape{Elem,Size}`).
- `tests/python/relax/test_vm_*.py`: 150 passed, 3 deselected
(`test_vm_multi_device.py` requires 3+ GPUs; host has 2 — env, not
regression),
2 xfailed.
- `tests/python/tirx-base/`: 273 passed, 2 skipped.
- `pre-commit` on edited files: green.
Closes#19443.
This PR fixes RPC tensor cleanup for tensors returned from remote calls.
When a remote function returns a `Tensor`, the RPC protocol sends both:
- the remote backing data pointer
- the remote tensor object handle used for deletion
Previously, `TensorFromRemoteOpaqueHandle` stored only the data pointer
and called
`FreeHandle(space_.data)` during local tensor destruction. That is
incorrect:
`FreeHandle` is meant for remote object handles, not raw data-space
pointers.
This could lead to invalid cleanup behavior and crashes during teardown
in RPC workflows, including the cross-compilation + RPC tutorial
scenario reported in #18923.
This change:
- stores the remote tensor object handle in `RemoteSpace`
- calls `FreeHandle(remote_tensor_handle)` during tensor destruction
- keeps cleanup fault-tolerant if the remote connection is already
closed
This PR brings up the tirx namespace. We have been spliting out the
original tir namespace to include high-level component s_tir and this PR
updates the remaining low-level part as tirx namespace
This PR phases out legacy target string format in favor of the json
style format that is more well formed. It also simplfies our overall
code in handling multiple formats.
This PR initalizes the s_tir for scheduable TensorIR. The change mainly
starts from python side, the we will gradually move towards the c++ side
in followup PRs. The python main change:
tir.Schedule => s_tir.Schedule
This PR renames tir.Block to SBlock. This clearly indicate the
scheduable property of the block and is a prereq for followup stir
passes refactor.
Main changes:
- Data structure change from Block to SBlock
- Syntax change from T.block to T.sblock
This PR cleans up the python API to make things more consistent
with existing python array api and torch.
Device update
- device_id => index, to be consistent with torch
- device_type => dlpack_device_type() returns int
- added type property same as torch.device
API updates:
- Move the convenient method like cpu() out into tvm runtime to keep device minimal
- tvm_ffi._init_api => tvm_ffi.init_ffi_api
- tvm_ffi.register_func => tvm_ffi.register_global_func
This PR Updates the NDArray => Tensor.
Both tensor and ndarray are commonly used terms.
Because the term Tensor is getting more common in the context of ML,
we do the rename to stay more aligned with torch.Tensor and DLTensor.
This PR formalizes original runtime::Module into ffi
as ffi.Module and cleans the APIs around it.
The goal is to stablize the Module API as extra API that can benefit the overall
ffi interactions. We also refactors the c++ code that depends on the Module.
[REFACTOR] Phase out getattr based attribute handling
This PR phases out getattar based attribute handling as they are slower
and introduces extra code path.
This does mean that if an Object is not explicitly registered
in python side, we will no longer be able to access the field by name.
Likely this is also desirable as we would like to enable faster use that
updates the python end and do not rely on these behavior.
This PR phases out tvm._ffi redirections in favor of new FFI
new functions are now called via tvm.ffi.
We also enabled limited API support for python 3.12+
so the compiled binary can be forward compatible to future
python versions.
This PR phases out the legacy c api in favor of the new FFI C API.
Also removes the redirection sccafolding for registry.h
- include <tvm/runtime/registry.h> => include <tvm/ffi/function.h>
- include <tvm/runtime/c_runtime_api.h> => include <tvm/runtime/base.h>
- TVM_REGISTER_GLOBAL => TVM_FFI_REGISTER_GLOBAL
The cleanup will greatly simplify the overall FFI surface of the project
and allows us to move towards an unified clean API based on tvm ffi.
This PR cleans up the container redirections and headers
so the files directly points to new ones in ffi folder
- runtime/shape_tuple.h => ffi/container/shape.h
- for IntTuple alias, introduce runtime/int_tuple.h for now
- runtime/container/array.h => ffi/container/array.h
- runtime/container/map.h => ffi/container/map.h
- runtime/container/optional.h => ffi/optional.h
- runtime/container/string.h => ffi/string.h
- runtime/container/variant.h => ffi/container/variant.h
- runtime/container/tuple.h => ffi/container/tuple.h
We also introduce limited number of tvm::ffi classes into tvm namespace,
when they are commonly used and their is no ambiguity.
This PR modernizes the FFI foundation of the project and introduce
a new minimal and lightweight module [tvm ffi](https://github.com/apache/tvm/tree/refactor-s3/ffi)
based on our lessons in the past few years. It implements a modern
version of the [Unified Packed and Object RFC](https://github.com/apache/tvm-rfcs/blob/main/rfcs/0097-unify-packed-and-object.md)
that unifies the packed function call and object systems.
Summary of the change:
- A dedicated clean Any/AnyView that can store strong and weak
references of items
- Function(previously PackedFunc) system built on top of the Any/AnyView
- A minimal C API that backs the overall calls. We are stabilizing the
API with a goal to bring clean, stable FFI conventions for both compiled
and registered code
- A rewrite of core python binding and generated code based on the module
- Update existing code and test cases to the new module
- Latest dlpack support
The new module brings many benefits thanks to the cleaner design,
to name a few:
- Any can support both POD types(int) and object types.
- Containers (e.g. Array) can now also contain Any value, e.g. now
`Array<int>` is supported, no need for boxed types
- Error handling now upgrades to object-based, allowing cleaner
traceback across languages
- Map now preserves insertion orders
- Path toward isolated stabilize minimum core ABI/API foundation module
- Type traits based design that cleanly defines how values interact
with Any system
- Automatic conversion of different types based on traits if needed
Because FFI upgrade is at heart of the project, the change touches every
component of the system. Importantly, this is an upgrade of the ABI so the
change is not backward compatible. The code compiled under the old
FFI won't work under the new one. We did provide example ABI translation
(e.g. LegacyTVMArgValueToFFIAny) functions for compatibility.
The PR tries to leave files in their old places while creating redirections.
The goal is to have the first milestone landed and infrastructure in place,
so we can do further refactors to complete features and cleanup legacy code
as trackable PRs. As of now, python binding and compiled code are under the
new convention while RPC and some other bindings still relies on legacy ABI
translation. We will work on upgrades in the coming PRs, including areas such
as reflection, phasing out legacy redirections etc.
* [REFACTOR] Phase out te.schedule python components
This PR phases out te.schedule python components.
te.compute is kept around for future usages.
tir.Schedule is a more modern version of the scheduling that we can use onwards.
Doing so also helps us to cleanup the testcases that relies on
explicit full build and execution. As we move future unit testcases
towards structural equality based unit tests.
* Simplify CI to focus on UT
The main rationale is that we should only have very few target
dependent UT in tests/python/codegen and possible
a new category in future for op-level integration if needed.
* Re-enable wasm
* fix lint
* remove hybrid,sparse autodoc and remove tests
---------
Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
This PR starts the step 0 to phase out relay from the current
development main branch. This PR focuses on the python
components of relay, autotvm, auto_scheduler. To make the change
manageable, we will also do followup steps on te.Schedule and
c++ components in followup PRs.
To continue support community members who depends on
legacy flows, the [v0.19.0](https://github.com/apache/tvm/tree/v0.19.0)
branch will continue contain these components.
As noted in [discussion on phasing out legacy components](https://discuss.tvm.apache.org/t/phasing-out-legacy-components/17703/30),
this would help us to do two purposes:
- By removing outdated or redundant elements, we can significantly
reduce complexity and improve maintainability.
- Unify our focus: Concentrating our efforts on the new unity flow
will allow for more efficient development and innovation.
It is also a good opportunity for us to revisit and reduce CI time.
The past relay legacy flow contains a lot of end to end tests that
requires hardware resources to run and causing long CI time.
Moving onwards, we can focus more on unit-tests that focuses
on structural equality and runs within seconds, while be mindful
about tests that requires hardware resources (by restricting them
to specific folders and CI nightly in some cases).
---
Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
The `PackedFunc` interface has arguments `int num_args` and `TVMValue*
args`, which contain the number of arguments and a pointer to the
array of arguments. Prior to this commit, when implementing the
`PackedFunc` interface for TIR `PrimFunc`s, the `MakePackedAPI` pass
would always assert that the `args` pointer was not null. However,
the `args` pointer is allowed to be null if `num_args` is zero. For
example, this occurs when calling an RPC function with no arguments.
This commit updates the `MakePackedAPI` transform to only assert that
`args` is non-null when `num_args` is greater than zero.
This PR enables RPCObjectRef over multi-hop RPC.
It is necessary to rewrap the argument as RPCObjectRef
so that the intermediate validation and re-encoding logic can
follow through.
[Runtime] Enable RPCObjectRef return in RPC
This PR enables RPCObjectRef return object similar to the disco transporation.
This allows us to do advanced remote debugging when remote vm requires
advanced object input like kv cache and shape.
To keep the implementation with minRPC(used in some of the limited protocols) forn now,
we only support RPCObjectRef for now and do not enable unpacking Shape and String.
The current unittest folder is too large and contains too many files and
too many components. This PR refactors the unittest folder by moving the
files to the corresponding folders.
Shushi Hong