Cleanup pass that relies on IR invariants instead of re-checking
already-guaranteed conditions. No new features; this is a
consolidation/cleanup pass only.
## Changes
- **docsifier (`python_doc_printer.cc`)**: the `ExprStringDoc` escape
scope always wraps the printer's fixed in-memory `ostringstream` sink,
which never short-writes and never enters a fail state. Drop the
streambuf-general short-write reporting in `xsputn`, the ctor `good()`
ICHECK, the dtor `rdstate`/`setstate` dance, and the redundant
post-render `good()` ICHECK; keep the one-line `saw_newline()` contract.
- **relax diagnostics (`well_formed.cc`, `block_builder.cc`)**: the ty
diagnostics test `ty.IsMissing()` on a now non-nullable `Type`, so word
them as "is missing" rather than "is nullptr".
- **relax numeric-gradient tests**: derive the device from the build
target via `tvm.device_from_target` inside the helpers instead of
threading a redundant `dev` argument that duplicates `target` at every
call site; annotate the numpy inputs as `np.ndarray`.
- **target/printer tests**: drop assertions that re-check a condition an
earlier assertion in the same test already guarantees.
## 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.
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 phases out the custom `tvm.testing.parameters()` helper in favor
of native `pytest.mark.parametrize`. `parameters()` itself is left in
place for now and removed in a follow-up, together with updating the
framework self-test
(`tests/python/testing/test_tvm_testing_features.py`) that exercises it.
Migration rules
- A group consumed only by test functions becomes
`pytest.mark.parametrize`.
- Single-name groups are unwrapped from 1-tuples to bare values.
- A group shared by multiple tests uses a module-level named list; a
test that uses only a subset of a group's names is parametrized only on
the names in its signature.
- `pytest.mark.parametrize` is stacked above the existing, unrelated
`tvm.testing.parametrize_targets(...)`, which is kept as-is.
Per-file pytest collection case counts are unchanged, except the two
intentional changes below.
Behavior changes (intentional)
- tests/python/relax/test_training_optimizer_numeric.py: the names `lr`
and `weight_decay` were rebound across three `parameters()` groups, so
`test_sgd` and `test_momentum_sgd` silently used the *adam* group's
`lr`/`weight_decay` (and `test_momentum_sgd` cross-producted with it:
2/6/2 = 10 cases). Native parametrize gives each test its own co-located
group: 2/3/2 = 7 cases. This fixes that latent rebinding bug; the case
count drops 10 -> 7 and `test_momentum_sgd` now exercises its own
`weight_decay` values.
- tests/python/target/test_arm_target.py: its `parameters()` group was
orphaned (no test consumed those names) — removed the dead definition.
Note: for tests that also use `tvm.testing.parametrize_targets`, the
generated test ids reorder the target (e.g. `test_unary[abs-True-llvm]`
-> `test_unary[llvm-abs-True]`); values and case counts are unchanged.
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.
## 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.
This updates CUDA fast math intrinsic lowering to use a PassContext
option instead of a CUDA Target attribute.
The new option is:
```python
with tvm.transform.PassContext(config={"tirx.enable_fast_math": True}):
...
```
When unset or false, CUDA math intrinsics continue to lower to the
precise CUDA math functions such as expf. When true, tirx.LowerIntrin
prioritizes the cuda.fastmath.* lowering rules, producing fast math
intrinsics such as __expf.
Fix CUDA lowering of standard TIR math intrinsics so they use precise
CUDA math functions by default instead of fast-math `__*f` functions.
This fixes the default behavior reported in #19546, where operators such
as `tirx.exp` could lower to `__expf` even though fast math was not
explicitly requested.
This change adds a CUDA target attribute, `enable_fast_math`, which
defaults to `false`. When the attribute is unset or false, standard math
intrinsics lower through the normal CUDA math rule, for example `expf`,
`logf`, `sinf`, `cosf`, `powf`, and `rsqrtf` for `float32`. When users
explicitly enable the attribute on the target, the lowering pass also
checks the `cuda.fastmath.FLowerIntrinsic` rules before the normal CUDA
lowering rules.
Users can opt in to fast math by constructing a CUDA target with the
attribute:
```py
tvm.target.Target({"kind": "cuda", "enable_fast_math": True})
target = tvm.target.Target({
"tag": "nvidia/nvidia-a100",
"enable_fast_math": True,
})
```
The fast-math lowering path currently covers the CUDA math operators
registered with `cuda.fastmath.FLowerIntrinsic`: `tirx.exp`,
`tirx.exp10`, `tirx.log`, `tirx.log2`, `tirx.log10`, `tirx.tan`,
`tirx.cos`, `tirx.sin`, `tirx.tanh`, and `tirx.pow`.
`tirx.rsqrt` is also registered for CUDA lowering so it maps to the CUDA
reciprocal-square-root intrinsic instead of being legalized as `1 /
sqrt(x)`.
Add CUDA codegen tests
`tests/python/codegen/test_target_codegen_cuda_fastmath.py` that check
the lowered IR, generated CUDA source, and runtime results for the
supported math intrinsics across floating point dtypes and both default
and fast-math targets.
## 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.
## Summary
This adds gating logic on top of #17699 to support optional subgroup
shuffle
primitives based on a compile-time flag.
## Problem
The PR #17699 always generates subgroup shuffle ops when targeting
WebGPU.
However, not all WebGPU devices support subgroups. We need a way to:
- Default to shared memory reductions (universally compatible)
- Optionally enable subgroup shuffles for devices that support them
## Solution
Implement gating via TVM target parameter:
- Default `thread_warp_size=1` disables warp reductions (uses shared
memory + barriers)
- Add target parser `UpdateWebGPUAttrs()` that sets
`thread_warp_size=32` when `supports_subgroups=true`
- Add `--enable-subgroups` CLI flag in mlc-llm to surface the option to
users
The gating happens at the reduction path selection level
(`IsWarpReduction()` in
`lower_thread_allreduce.cc`), ensuring subgroup ops are never generated
unless explicitly enabled.
## Testing
Tested with Llama-3.2-1B-q4f16_1. Baseline (no flag) uses shared memory
reductions;
with flag, generates subgroupShuffle* ops.
Both the generated WGSLs here:
https://gist.github.com/ksgr5566/301664a5dda3e46f44092be4d09b2d4f
Benchmarking:
https://gist.github.com/ksgr5566/c9bd5bc5aadba999ec2f2c38eb0c49b3
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 cleans up the target python api.
- Removes the indirections of attribute exposure
- Move tag registry to python so it is easily configurable
- Remove legacy constructors in favor of tags
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 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 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.
* cleanup relay c++
* [REFACTOR] Phase out relay c++ components
This PR phases out the relay C++ components and
simplifies the overall codegen runtime logic.
---------
Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
* [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>
This PR phases out VTA from the current development main branch.
The particular component will remain available in past releases
and is not actively maintained as of now.
Currently, target features are determined by a set of fixed checks on
the target string. This works well for checking support of a small
number of simple features, but it doesn't scale. Some problems include:
- There are many non-trivial conditions for which a feature may(not) be
available. It is easy to miss these with the current implementation.
- The inclusion of some features in a target string can imply other
features. For example, "+sve" implies "+neon". This currently isn't
taken into account.
- The tests in tests/cpp/target/parsers/aprofile_test.c suggest that
targets such as "llvm -mcpu=cortex-a+neon" and "llvm -mattr=+noneon"
are supported target strings. The features will be correctly parsed in
TVM, however, they are not valid in LLVM. Therefore, it's possible
that TVM and LLVM have different understanding of the features
available.
This commit uses the more robust LLVM target parser to determine support
for the features in TVM. It leverages previous infrastructure added to
TVM for obtaining a list of all supported features given an input
target, and uses this to check the existance of certain features we're
interested in. It should be trivial to grow this list over time. As a
result of this change, the problems mentioned above are solved.
In the current form, this commit drops support for target strings such
as "llvm -mcpu=cortex-a+neon" and "llvm -mattr=+noneon". A scan of the
codebase suggests this functionality is not in use (only in test cases).
Should we feel the need to support them, or have a smoother migration
for downstream users of TVM we can add a translator to the parser to
convert these into LLVM compatible targets.
This commit adds support for generating code for scalable loads and
stores. It also adds support for the creation of scalable broadcast
operations.
Co-authored-by: Elen Kalda <elen.kalda@arm.com>
Co-authored-by: Neil Hickey <neil.hickey@arm.com>
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.
This PR leverage LLVM itself for CPU features lookup, replacing hard-coded lists.
In order to keep maintainability with X86 families & features we can rely on LLVM.
---
Changes:
* Introduce a single ```target_has_feature(XXX)``` replacing all ```target_has_XXX()```
* PY+FFI: expose new ```llvm_x86_get_archlist```, ```llvm_x86_get_features``` & ```llvm_x86_has_feature```
* PY: expose new ```target_has_feature``` wrapper to ```_ffi.llvm_x86_has_feature```
---
There is a test unit for a comprehensive check with the old behaviour.
For better reliability, this way of feature checking can be implemented for other arches.
* Add Python representation for VirtualDevice
This adds a Python class to represent the VirtualDevice so that the
behaviour for `device_type()` can be semi-replicated.
These tests were actually not being ran and were broken so I've added
them to the integration script.
* Update other references to make_virtual_device
[Target] Adds SEScope (Storage/Execution Scope) for use as new unit of planning in 'device' planning
This is the first step in https://github.com/apache/tvm-rfcs/pull/38 to bring devices
and targets together when doing device planning. I've gone ahead and also included a
memory scope in this object since we will also need to propagate memory scopes across
Relay expressions once this basic preparation is in place. In the meantime that field will be
left as "".
Once device planning works in units of SEScopes it will be possible to directly read off
the device and target for any Relay sub-expression without the need for TargetMaps ort
the construction of default Targets.
SEScopes also support 'Join' and 'Default' operations needed when constraint solving in
the device planner. You can see those in use in my scratchpad branch:
https://github.com/mbs-octoml/mbs-tvm/tree/mbs-scopes
This PR also brings some duplicated and the ad-hoc 'default target' handling logic
together into a CompilationConfig class. (Again, see the scratchpad branch for how that
will end up being used). I've placed that next to SEScope since it's main purpose is to
a) establish the default SEScope for primitive ops
b) establish the SEScope for the 'host'
c) feed a definitive vector of Targets into device planning so it can resolve all
"on_device" and "device_copy" device references to their full SEScope form.
* Reworked to avoid global SEScopeCache.
Realized while working through unit tests in the sequel that it's reasonable
for folks to call build multiple times with distinct Target objects, in which
case the global cache would grow without bound.
So instead placed the cache in the CompilationConfig class. Since that class
now has everything the device planner needs to do its job, promoted it to
be an FFI-able Object, which is now in compilation_config.{h,cc}.
I think we can do much better with CompilationConfig, but for now keeping it
to the minimum I needed to prepare for device planning from all the executor
compilation codepaths.
Tianqi Chen