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.
In the past we have been using `DataType` in PrimExpr.dtype field to
check type information for PrimExpr while still having BaseExpr.ty for
richer type information. DataType is also used both in runtime and
compiler. This PR streamlines the boundary:
- PrimExpr.ty now carries PrimType that replaces original use of
`DataType`
- Runtime use will now favor DLPack DLDataType, removing one layer of
indirection.
- Constants attributes where values are usually runtime values, will use
`DLDataType`
- DataType will be phased out after this PR
We also brings up helper functions in PrimType, but also limits them to
a more concise set so the functions do not grow with the data type codes
in DLPack.
This is a major refactor that changes the IR primitive. It helps to
bring possible future benefits:
- Unified type mechanism through Expr.ty
- Possibility of carry future Type nodes
Migration Guide:
- Use `PrimType` when code reasons about compiler expression types,
tensor element compiler types, or constructs a `PrimExpr`/compiler type.
- Use existing source types such as `expr.ty()`, `ExprOp.expr_ty()`, or
TE tensor element `dtype` where possible instead of rebuilding a type
from dtype text.
- Use raw `DLDataType` for runtime constants, ABI paths, dtype-valued
attrs, and storage/runtime helper logic.
- Prefer direct `PrimType` equality, `MatchesCode(...)`,
`MatchesElementType(...)`, and `WithCode(...)` over local wrappers or
string dtype checks.
Performance:
Using Object type instead of DLDataType would indeed bring some
performance impact to the IR. We have done the following performance
optimizations:
- Make sure most of the outputs reuse one of the PrimType from inputs
- Cache a thread local PrimType based on input so we don't repeatly
realloc
We did benchmarks show that rewrite simplify operation stays within
+-10% overhead of original one. Which merits the refactor given the
benefit the unfication brings
## Summary
- unify Relax's former StructInfo surface into the Type vocabulary and
Expr.ty storage path
- remove leftover DependentTypeNode and legacy OpNode::op_type storage
- keep base Type nullable while concrete Relax/DTensor type refs are
non-nullable
- clean stale StructInfo/TensorStructInfo/sinfo vocabulary in
Python/docs and distributed-op macros
- address Gemini follow-ups for parser annotations, BlockBuilder
docstring, and Adreno TensorType cast audit
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.
## Summary
The old Hexagon app and test wrappers depend on RPC helper artifacts
that are no longer part of the supported app flow. This PR removes those
wrappers and related helper references while keeping the core Hexagon
target, codegen, and runtime implementation in place.
## Changes
- Remove the obsolete Hexagon app wrapper directories.
- Remove the Hexagon contrib test directory and its dedicated pytest/RPC
launcher helpers.
- Drop stale CI/docs references to the removed app and test helper
paths.
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`
#19714 intentionally removed packed sub-byte (bits 1/4) support from
RandomEngine after the int4 fill path was found to write past the
allocated storage, but left "int4" in test_random_fill's dtype list, so
the test now fails with InternalError on every run. Remove it from the
list to match the supported set.
Also cast dtype.code / dtype.bits to int in the unsupported-dtype error
messages: both fields are uint8_t and were streamed as (invisible)
characters, e.g. "Doesn't support dtype code dtype bits".
The test_hexagon directory is not enrolled in CI and has drifted out of
sync with the codebase:
- Delete test_fixed_point_conversion.py: it imports
tvm.topi.hexagon.utils, but the tvm.topi.hexagon module (including
get_fixed_point_value under test) no longer exists in the source tree.
- Delete test_2d_physical_buffers.py: its ir_module fixture was removed
in #17665 (Feb 2025), leaving all 36 tests erroring at setup with
"fixture 'ir_module' not found" ever since.
- Fix test_take.py: drop the stale exec_mode="compiled" argument that
tvm.compile no longer accepts. Both tests in the file pass again (llvm
host target, no Hexagon hardware required).
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.
Previously `_generate_triton_kernel` overwrote the user-provided kwargs
with the constexpr dict before calling triton.compiler.compile, so
options such as num_warps passed to T.call_kernel were silently dropped.
Pass constexprs to ASTSource and forward the user kwargs as compile
options.
The pre-3.3 compatibility branches are removed in favor of an explicit
minimum-version check: they were never exercised in CI (which does not
install Triton), and Triton >= 3.3 has shipped with PyTorch since 2.7.
The integration test now matches the actual lowering, where constexpr
parameters (BLOCK_SIZE) appear as runtime kernel arguments in
call_packed, and passes num_warps=8 expecting a thread extent of 256 to
cover the option forwarding.
This pr fixes two contrib test failures that show up when running
source-tree tests against the `apache-tvm` wheel.
- Launch the pickle memoize helper script with `sys.executable` so
subprocesses use the same Python environment as pytest.
- Fix `tvm.contrib.random.random_fill` for packed sub-byte dtypes by
using the actual tensor storage size instead of the logical element
count.
## Root Cause
`tests/python/contrib/test_memoize.py` executed
`pickle_memoize_script.py` directly. In a wheel test environment, the
shebang can resolve to a different Python than the active wheel venv,
causing `ModuleNotFoundError: No module named 'tvm'`.
`random_fill` used the product of tensor shape as the number of values
to write. For packed dtypes such as `int4`, two logical elements share
one byte, so this wrote past the allocated storage and caused native
heap corruption / process abort.
---------
Co-authored-by: tqchen <tianqi.tchen@gmail.com>
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
Python callers should reach the canonical tvm-ffi structural helpers
directly instead of going through a TVM-side redirect layer. This makes
the public tvm.ir bindings exact aliases of the tvm_ffi APIs and exposes
get_first_structural_mismatch from tvm.ir.
Main changes:
- Import structural_equal, get_first_structural_mismatch, and
structural_hash directly from tvm_ffi
- Remove the pure wrappers from tvm.ir.base while keeping
assert_structural_equal's TVM-specific formatting
- Update mismatch tests and add identity coverage for the direct
bindings
In continuation of #19624 this catches some unlifted entries.
Hope there is no more left, for consistency it now covers comments and
perhaps non-active (hotpath) parts.
## 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 PR adds annotation support to `tirx.Call` so downstream codegen
users can attach call-level metadata and preserve it through TIRX
transforms.
What changed:
- Add `CallNode::annotations` and expose it through reflection.
- Add Python `tvm.tirx.Call(..., annotations=...)` support.
- Preserve call annotations in C++ and Python expression mutators.
- Preserve annotations across TIRX/arith passes that rebuild equivalent
calls.
- Print annotated calls as `Tx.Call(..., annotations={...})` and support
script roundtrip.
- Add regression coverage for annotated calls, mutator preservation,
script roundtrip, and simplify preservation.
This pr also cleans some stuff that #19596 didn't clean completely
Fixes #<issue-number>.
Reads of `_msg_size` from the tracker socket are now bounded to
`MAX_TRACKER_MSG_BYTES = 1 MiB`, and the 4-byte size header is
consumed at read time. Without these checks, a single TCP connection
from a peer can grow the tracker process buffer until OOM, and a wire
size of 0 starves the parser without ever freeing the bytes.
Per the TVM security model the tracker is deployed on trusted networks,
so this is filed as a robustness defect, not a security advisory.
Apache security team triage (private thread, 2026-05-17) confirmed this
is the right channel.
### Test
Added regression test in tests/python/contrib/test_rpc_tracker.py that
completes the magic handshake, sends an oversized msg_size header
(0x7FFFFFFF), and asserts the tracker closes the connection.
### Changes
- python/tvm/rpc/tracker.py: bound `_msg_size` to (0,
MAX_TRACKER_MSG_BYTES], consume size header on read.
- tests/python/contrib/test_rpc_tracker.py: regression test.
## 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.
Supersedes #18247. Per maintainer guidance, resubmitting as a fresh PR
due to CI workflow changes affecting old PRs.
## Summary
This PR adds an example NPU BYOC backend for Relax, including end-to-end
integration points:
- pattern registration
(`python/tvm/relax/backend/contrib/example_npu/patterns.py`)
- backend registration
(`python/tvm/relax/backend/contrib/example_npu/__init__.py`)
- codegen entrypoint
(`src/relax/backend/contrib/example_npu/codegen.cc`)
- runtime module
(`src/runtime/contrib/example_npu/example_npu_runtime.cc`)
- CMake integration (`cmake/modules/contrib/ExampleNPU.cmake`,
`CMakeLists.txt`, `cmake/modules/LibInfo.cmake`,
`src/support/libinfo.cc`)
- tutorial/docs (`docs/how_to/tutorials/byoc_npu_example.py`, README
under contrib path)
- tests (`tests/python/contrib/test_example_npu.py`)
- CI build config enablement (`tests/scripts/task_config_build_cpu.sh`)
## Review feedback addressed from #18247
- Test location under `tests/python/contrib/`
- README includes explicit enable instructions for
`USE_EXAMPLE_NPU_CODEGEN` and `USE_EXAMPLE_NPU_RUNTIME`
- README quick-start uses inline `MatmulReLU` (no import from test
module)
- Added CMake source wiring and feature flags for runtime/codegen
- Added docs tutorial under `docs/how_to/tutorials/` (not only README)
- Reorganized motivation/context section near top of README
- Extended pattern coverage to include `example_npu.softmax` with tests
## Validation
Local checks run:
- `pre-commit` on touched files (pass)
- `PYTHONPATH=python python -m pytest -q
tests/python/contrib/test_example_npu.py` (pass)
## Notes
This backend is an example/tutorial implementation (CPU-emulated)
intended to document modern NPU-oriented BYOC integration patterns and
provide a reference path for future hardware-specific backends.
## Summary
`pickle_memoize` loads cached pickle files via `pickle.load()` without
any integrity verification or user warning. If an attacker can write to
the cache directory, they can inject malicious pickle payloads that
execute arbitrary code on next load.
## Fix
Adds a `UserWarning` when loading pickle cache files to alert users
about the security risk.
## Related
Huntr security vulnerability report (CWE-502: Deserialization of
Untrusted Data)
Signed-off-by: scruge1 <scruge1@proton.me>
Co-authored-by: scruge1 <scruge1@proton.me>
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
## Summary
This PR introduces `AllocBufferNode`/`AllocBuffer` as a single TIR
statement that both allocates memory and declares a buffer into scope.
This replaces the previous pattern of `Allocate(var, dtype, shape, cond,
DeclBuffer(buf, body))` with the simpler `AllocBuffer(buf, body)`.
### Main changes
- **New IR node** `AllocBufferNode` with fields `{buffer, annotations,
body}` — same semantics as `DeclBuffer` but also allocates memory
- **TVMScript**: `T.alloc_buffer(shape, dtype, scope)` now emits
`AllocBuffer` directly (statement-level allocation).
`T.sblock_alloc_buffer(...)` for SBlock-level buffer allocation (full
parameter set)
- **All codegen backends** (C, CUDA, Metal, OpenCL, WebGPU, LLVM, NVPTX,
AMDGPU, SPIR-V) updated to handle `AllocBufferNode`
- **All TIR transforms** (storage_rewrite, flatten_buffer,
vectorize_loop, lower_warp_memory, etc.) updated
- **All S-TIR transforms** (compact_buffer_region, merge_shared_memory,
inject_double_buffer, etc.) updated
- **Removed `AllocateNode`** entirely — `AllocBuffer` is now the sole
allocation primitive
- **Removed `AllocDescriptor`** from merge_shared_memory_allocations —
uses `Buffer` objects directly
- **Added `AllocBuffer::ConstantAllocationSize()`** inline helper method
### Design rationale
The old `Allocate + DeclBuffer` pair was a historical artifact:
`AllocateNode` stored raw fields (`buffer_var`, `dtype`, `extents`,
`condition`) separate from the `Buffer` object, requiring pattern
matching (`IsAllocateDeclBufferPattern`) to reconstruct the buffer
association. `AllocBuffer` unifies this into a single node with a proper
`Buffer` reference, simplifying codegen backends and transform passes.
225 files changed, ~3500 insertions/deletions (net near-zero, mostly
mechanical migration).
## Test plan
- [x] All TIR base tests pass
- [x] All TIR transform tests pass
- [x] TVMScript roundtrip tests pass
- [x] S-TIR transform tests pass
- [x] Codegen tests pass
- [x] All-platform minimal tests pass
- [x] C++ functor tests pass
- [x] Pre-commit clean (clang-format, ruff, etc.)
## Summary
Remove 9 test scripts from `tests/scripts/` that are no longer used in
CI —
i386 and hexagon architectures were removed from CI, mypy moved to
pre-commit,
and several others are unreferenced.
This PR enables ruff pyupgrade (UP) rules with py310 target, auto-fixing
~5600 annotation modernizations (PEP 585 generics, PEP 604 unions,
deprecated typing imports).
Also removes from __future__ import annotations from ir/module.py and
rmsnorm.py, bumps requires-python to >=3.10, and removes absolute_import
aliases from topi/contrib files.
## What
Replace 6 bare `except:` clauses with `except Exception:`.
## Why
Bare `except:` catches `BaseException`, including `KeyboardInterrupt`
and `SystemExit`, which can prevent clean process shutdown and mask
critical errors. Using `except Exception:` catches all application-level
errors while allowing system-level exceptions to propagate correctly.
This PR removes legacy runtime contrib backends that have no existing
compiler backend,
no active development. They can always be brought back in future in case
we find there is a need
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 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
## Why
The scatter_elements operator was missing FRelaxInferLayout support,
which prevented proper layout transformation when used with operators
like conv2d that require layout conversion.
## How
- Implement InferLayoutScatterElements function that handles layout
inference for scatter_elements
- Transform axis attribute according to the inferred layout using
FindAxis
- Handle sub-indexed layout fallback to initial layout
- Add test case for conv2d + scatter_elements layout conversion
This PR updates TVM_FFI_STATIC_INIT_BLOCK to function style.
Now we do the code as follows, which is cleaner in generally
and also helps error reporting to locate the right place.
```
TVM_FFI_STATIC_INIT_BLOCK() {
RegisterStaffs();
}
```
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.
Previously we simply take the raw symbol for DSO libraries.
This can cause symbol conflict of functions that take the ffi calling convention
and those that are not.
This PR updates the convention to ask for LLVM and libary module to always
append a prefix __tvm_ffi_ to function symbols, this way we will no longer
have conflict in TVM_FFI_EXPORT_DLL_TYPED macro
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.
This PR refactors the tir.call_llvm_intrin to omit the first nargs argument in the beginning.
Previously the nargs was introduced when prefetch have different number of signature.
The previous reason no longer stands as of now, and it is less intuitive to attach nargs
for the call_llvm_intrin, where nargs directly appears in number of argument.
After the update, tir.call_llvm_intrin can directly pass in the arguments as it is.
Tianqi Chen