## Summary
Compiler Targets can carry device-type semantics that runtime
device-name parsing does not preserve.
- add `tvm.device_from_target` for canonical Target-to-Device
translation
- use explicit runtime constructors where the device kind is fixed
- update target-derived utilities, tests, and documentation to use the
explicit boundary
Add tvm.testing.run_with_gpu_lock backed by the existing
tvm_ffi.utils.FileLock. Migrate live local GPU tests to acquire the
machine-local lock around device execution, synchronization, host
transfer, and checks while leaving target construction and compilation
outside the critical section.
Replace the custom xdist scheduler with standard xdist_group placement
for the order-dependent test family. RPC tests retain dynamic port
allocation and per-test process isolation rather than gaining a broad
category lock.
## 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.
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 is the Follow-up to #19777. This pr removes the last
`requires_*` decorators so test gating is plain pytest everywhere, with
no custom indirection left.
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`
The CUDA-specific timer preprocessor test was the only user of the
nvbench-derived L2 cache flush helper. Removing that path avoids
carrying a special CUDA registration and third-party header for
benchmark-only test coverage.
This PR removes the CUDA L2 cache flush registration, the dedicated
preprocessor timing test, and the now-unused nvbench-derived
license/dependency entry.
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 refactors and establishes ffi.Module under the python tvm ffi api.
Also moves export_library method to executable so it aligns more with
compiled artifact.
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.
* [FFI] Phase out ctypes fallback in favor of cython
In the past we maintained both ctypes fallback and cython path for our FFI.
One reason was cython was not that widely available then.
The situation has changed, so this PR removes the ctypes fallback
so we can focus on building cython based core for FFI.
Notably, the cython core still leverages ctypes, especially for
some of the features that do not need fast path, and primarily
uses cython to speedup core FFI features with fast path.
With cython, we can get most FFI costs down to level
of roughtly 1e-6 sec, which aligns with most of the costs like
overhead of a numpy operation.
* fix lint
* Fix the dso loading issue
This refactor introduces a base Executable class and a `tvm.compile`
interface that can be used to compile both TIR and Relax programs.
`tvm.compile` will return an Executable object that can be used to call
either TIR or Relax functions.
* [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>
* Revert "Revert "[FFI][RUNTIME] Introduce runtime boxed types for int/float/bool" (#17252)"
This reverts commit 11be832620.
* [FFI] Re-introduce the boxed primitive values
Initially introduced in https://github.com/apache/tvm/pull/16183,
these changes were reverted in
https://github.com/apache/tvm/pull/17252 due to performance
degredation in some Relax models. This could occur when a model
contained a large number of calls to `"vm.builtin.tuple_getitem"`,
which may occur when model weights are provided as a tuple.
This PR re-applies the changes from
https://github.com/apache/tvm/pull/16183, but with the performance
degredation resolved. The root cause was unnecessary type-checking
when converting from an untyped `tvm::ArrayNode*` to the typed
`tvm::Array<T>`, in the case where `T` is `ObjectRef`.
* Correct typo from T to U
* [Container] Support non-nullable types in Array::Map
Prior to this commit, the `Array::Map` member function could only be
applied to nullable object types. This was due to the internal use of
`U()` as the default value for initializing the output `ArrayNode`, where
`U` is the return type of the mapping function. This default
constructor is only available for nullable types, and would result in
a compile-time failure for non-nullable types.
This commit replaces `U()` with `ObjectRef()` in `Array::Map`,
removing this limitation. Since all items in the output array are
overwritten before returning to the calling scope, initializing the
output array with `ObjectRef()` does not violate type safety.
* [FFI] Separate runtime types from IR types for int/float/bool
Prior to this commit, `int`, `float`, and `bool` arguments from Python
were converted to `IntImm`, `FloatImm`, and `Bool`. These are
subtypes of `PrimExpr`, and should only be used at compile-time. By
automatically applying this conversion as part of the FFI, these types
are required to be present whenever a primitive is converted to a
`tvm::ObjectRef`.
This can become especially fragile for an end-user when storing
objects into a TVM container. Because TVM containers require all
contents to be `ObjectRef` subclasses, an automatic conversion may be
applied on storing into a container, resulting in an unexpected type
being retrieved from the container. For example, this currently
occurs in Relax when extracting a `R.Prim` from a `R.Tuple`.
This commit introduces a `Box<T>` type for storage of boxed primitives
at runtime, distinct from the IR types.
* Primitive arguments provided to a PackedFunc that requires an
`ObjectRef` will be converted to the corresponding boxed type.
(e.g. Passing a Python `int` to a C++ function accepting `ObjectRef`
produces a `Box<int64_t>`.
* Boxed primitives provided to a PackedFunc that requires an unboxed
primitive will be converted to the corresponding primitive.
* PackedFunc return values of `ObjectRef` are converted to the
corresponding primitive, if present. (e.g. If a `tuple_getitem`
with static return type `ObjectRef` returns a `Box<int64_t>`, it
will be unwrapped to a python `int`.)
Together, these three rules provide backwards compatibility for
existing PackedFunc definitions, while avoiding exposing the user to
any container-induced type conversions betweeen primitive types and
`ObjectRef`.
* Fix unit test failure after merge
* Fix breakage in new unit test
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.
* [Runtime] Allow offset to be specified in NDArray::CreateView
Prior to this commit, the `NDArray::CreateView` method could produce
an aliasing view of an existing array with a different shape or
datatype, but the view was required to have the same
`DLTensor::byte_offset` as the existing array. This commit updates
the `NDArray::CreateView` method with an additional parameter,
specifying the offset of the view relative to the existing array.
* Change type of `relative_byte_offset` from `size_t` to `uint64_t`
Both to match the type used in `DLTensor::byte_offset`, and to resolve
compilation errors on 32-bit platforms, which fail to compile due to a
missing `Type2Str` specialization.
This commit resolves a bug that was introduced in
https://github.com/apache/tvm/pull/16377. If no CUDA-capable GPUs are
present, the call to `cudaGetDeviceCount` will return an error, which will
be raised as an exception by the `CUDA_CALL` macro. However, checking
the `kExist` flag is valid even if no GPUs are present.
This commit removes the use of `CUDA_CALL`, and instead returns false
in this case.
Currently, when a default compile target such as llvm is specified,
it implies llvm -keys=cpu which tends to imply x86 related components
being used during compilation e.g. the schedules registered in TOPI.
This can be confusing for a user when compiling on other architectures,
especially when other tools such as llc infer the default target
based on the host.
When the target kind is llvm, this commit uses the
"target.llvm_get_system_triple" functionality to automatically detect
mtriple when one has not been provided in the target string. The
target will be updated to one that uses the mtriple of the host:
llvm -> llvm -mtriple=<system-triple>. When compiling on Arm(R)-based
targets, this has the added benfit of automatially introducing
-keys=arm_cpu to the target improving the schedule selection.
Lots of tests are currently using targets such as llvm or similar
which has resulted in a lack of coverage of other targets such as
arm_cpu. As part of this commit, failing test cases which have simple
/ obvious issues have been fixed. Others that likely need more thought
have been skipped. In doing so, it reduces the number of modifications
and simplifies the review for this change.
This commit is a follow up of the changes made in: #14981
Change-Id: Icee7f5c00d58fc77367c823273fccae128260471
Co-authored-by: Jack Frankland <jack.frankland@arm.com>
---------
Co-authored-by: Jack Frankland <jack.frankland@arm.com>
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.
Tianqi Chen