103 Commits

Author SHA1 Message Date
Tianqi Chen e479a5dbe7 [RUNTIME][PYTHON] Add explicit Target device conversion (#20005)
## 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
2026-07-15 05:34:21 +08:00
Tianqi Chen bbfdab79d9 [CI] Repair Python test cleanup regressions (#19955)
## Summary

- Keep the Python test launcher close to plain `pytest -n auto`, move
nightly tests under `tests/nightly/python`, remove obsolete launchers
and collection bookkeeping, and partition CPU/GPU jobs with explicit
`gpu` marker expressions.
- Repair exact-pointer regressions at their owning boundaries: packed
raw-string ABI values, CUDA/Metal matrix intrinsic pointers, internal TE
extern offsets, MetaSchedule scalar annotations, localized
auto-tensorization scope matching, and typed DLTensor fixture fields.
- Preserve typed workspace calls in TIR and cast pointer-returning
external calls in CodeGenC, covered by a plain-TIRx 1024-byte global
workspace that is compiled as C++.
- Finish phasing out value-bearing Relax `R.Prim` annotations by
requiring an explicit dtype, removing obsolete value-based contracts,
and expressing the DISCO rank-dependent slices as explicit scalar
`call_tir` inputs.
- Gate the distributed callback on the optional DISCO runtime, NCCL, and
at least two GPUs so capability-limited jobs skip instead of failing.
- Remove the non-demonstrating pointer probe, use direct TVMScript
comparison for packed strings, and remove the four designated legacy
testing modules.

The seven repaired CPU categories cover packed raw strings (7 failures),
CUDA/Metal matrix access-pointer types (7), internal TE extern offsets
(1), a typed DLTensor fixture (1), MetaSchedule scalar annotations (1),
CodeGenC workspace return casts (12), and localized auto-tensorization
storage-scope matching (19).

## Validation

- Base: `ded6ad8dd212869c881efb5590f8a33fc972728e`
- Head: `a7277e86dbcfe0638c8c252d36760859c4ab4297`
- All 35 locally available original failing node IDs pass across the
focused runs.
- The full focused TE, TIR builtin-lowering, and CodeGenC files pass: 61
tests.
- The complete touched Relax/TVMScript set plus
PlanAndUpdateBufferAllocationLocation passes with 784 passed, 20
skipped, and 1 expected failure.
- The DISCO callback collects and skips when its runtime or two-GPU
environment is unavailable.
- Six direct mapping tests, twelve tensor-core sketches, and the dp4a
sketch pass unchanged.
- The compiler rebuild, branch-wide pre-commit hooks, and full-range
whitespace checks pass.
- The 13 broad CBLAS/TFLite nodes remain dependency-gated; their owning
TE and generated-C regressions compile.

No merge is included in this change.
2026-07-06 16:29:52 +08:00
Tianqi Chen 275114b327 [REFACTOR][IR] Unify PrimExpr with Expr typed view (#19910)
## Summary
- Make `PrimExpr` a typed C++ view over `Expr` values whose
`ExprNode::ty` is `PrimType`, instead of using a separate runtime node
class as the proof of primitive-ness.
- Use the shared `ir::Call` node for Relax, TIRX, and primitive-valued
calls, while keeping primitive-only APIs explicit at their semantic
boundaries.
- Keep Python on the general `Expr` surface for primitive-typed values
so `isinstance` behavior does not imply a nominal primitive-expression
subclass.

## Design Rationale
The main advantage of this change is that common expression nodes such
as `Call` can be unified without specializing each one to `PrimType`. A
single `ir::Call` can represent a Relax tensor call, a Relax scalar
call, or a primitive-valued intrinsic call; the result type stored in
`ExprNode::ty` determines whether that particular value can be viewed as
`PrimExpr`.

This keeps the IR node hierarchy focused on expression structure rather
than result-type categories. Nodes that are intrinsically primitive,
such as integer and floating-point literals or TIRX primitive operators,
still have strongly typed C++ APIs and data structures. General nodes
whose result type may vary, such as `Call`, remain general `Expr` nodes
and are narrowed to `PrimExpr` only where primitive-only semantics are
required.

The PR also keeps the compatibility surface practical: C++
primitive-only APIs continue to accept `PrimExpr`, Python exposes a
compatibility predicate for checking the primitive typed category, and
visitors/printers use one natural `Call` path rather than duplicating
Relax and primitive call handling. Missing expression types are
represented explicitly with `Type::Missing()` so constructors can leave
type inference to later analysis without relying on nullable `Type`
values.
2026-07-01 18:55:33 -04:00
Shushi Hong f2a584a5b8 [Tests] Remove dead helpers and unused probes from tvm.testing (#19821)
Drop accumulated dead code in the test-support package: helpers with
zero call sites, unused capability probes, dead FFI re-exports, and
orphaned pytest plumbing. Verified by repo-wide grep that nothing
references any of these.
2026-06-18 16:14:31 -04:00
Shushi Hong a8635b04cc [Tests] Migrate off tvm.testing.parametrize_targets to native pytest (#19826)
This pr moves target selection and per-target device gating off the TVM
pytest plugin onto plain pytest, and remove the now-dead machinery.
2026-06-18 08:48:53 -04:00
Shushi Hong ddfec9c3d6 [Tests] Remove the now-unused tvm.testing.parameters() helper (#19807)
All in-tree uses of tvm.testing.parameters() were migrated to native
pytest.mark.parametrize in #19803, so remove the helper itself along
with the plugin machinery that only served it:
- python/tvm/testing/utils.py: delete the parameters() function and the
_parametrize_group counter.
- python/tvm/testing/plugin.py: delete
_parametrize_correlated_parameters and its call in
pytest_generate_tests.
- tests/python/testing/test_tvm_testing_features.py: drop the
joint-parameter tests that exercised parameters() (the parameter() and
fixture() tests stay).

This removes the public tvm.testing.parameters symbol;
tvm.testing.parameter (singular) and tvm.testing.fixture are unchanged.
Use pytest.mark.parametrize instead.
2026-06-17 00:32:33 -04:00
Shushi Hong 35a35b8434 [Tests][Refactor] Remove unused testing helpers (#19800)
CompareBeforeAfter, skip_parameterizations, and xfail_parameterizations
have no remaining users anywhere in the repo. CompareBeforeAfter (a base
class for TIR before/after transform tests) has been superseded by the
inline assert_structural_equal(transform(Before), Expected) pattern, and
the {skip,xfail}_parameterizations helpers (which marked specific
parametrizations at runtime) are unused -- native pytest.param(...,
marks=...) covers that need.

Also drop the private _mark_parameterizations helper they relied on and
the now-unused 'import textwrap'.
2026-06-16 17:51:23 -04:00
Shushi Hong 949be814c0 [Docs] Modernize test-gating documentation (#19788)
This pr updates the contributor guide and tvm.testing
docstrings/comments to describe the current gating API

---------

Co-authored-by: Tianqi Chen <tqchen@users.noreply.github.com>
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
2026-06-16 21:09:31 +08:00
Shushi Hong 9011739dc2 [Tests] Replace remaining requires_* helpers with standard pytest (#19787)
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.
2026-06-16 07:32:37 -04:00
Shushi Hong e4da848e57 [Tests] Modernize test gating (#19777)
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`
2026-06-15 18:50:57 -04:00
Shushi Hong b57d0b3d07 [Runtime][Disco] Fix session attribute storage, NVSHMEM build, and test gating (#19736)
The tvm_ffi Object metaclass now gives every subclass `__slots__ = ()`,
so the Disco Python wrappers can no longer store instance attributes and
every session construction fails with AttributeError. Declare the
attributes each
wrapper actually stores as named slots, fix the NVSHMEM `dist_gemm.cu`
so TVM builds with `USE_NVSHMEM = ON`, and gate the disco tests on the
disco runtime being present so they skip cleanly on builds (e.g. the pip
wheel) that report `USE_NCCL` / `USE_NVSHMEM = ON` without shipping it.

### Session attribute storage
- `DPackedFunc` / `DModule`: `__slots__ = ("session",)`.
- `Session`: `__slots__ = ("_cache", "_import_python_module")`
2026-06-12 08:32:23 -04:00
Tianqi Chen 96cba60464 [PYTHON] Autoload backends; simplify library loading; remove TVMError for native errors (#19727)
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.
2026-06-11 13:50:38 -04:00
Tianqi Chen ffea531107 [REFACTOR][PYTHON] Lift compiler/CLI/process modules from tvm.contrib to tvm.support (#19624)
## 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.
2026-05-27 15:31:12 -04:00
Tianqi Chen 02b130249c [REFACTOR][TIR][ARITH] Phase out ControlFlowGraph, NarrowPredicateExpression, and rename Simplify to StmtSimplify (#19604)
## Summary

This PR cleans up technical debt in the TIR simplification machinery via
two commits:

**Commit 1: Phase out ControlFlowGraph and NarrowPredicateExpression**

- Remove `ControlFlowGraph` (~2360 lines) from `src/tirx/analysis/` —
used only in
  non-default config paths that are no longer maintained
- Remove `NarrowPredicateExpression` from `src/arith/` — sole non-test
caller was `ControlFlowGraph`
- Remove gated config fields `propagate_knowns_to_prove_conditional` and
  `propagate_knowns_to_simplify_expressions` from `SimplifyConfig`
- Remove `use_dataflow_analysis` from `RemoveNoOpConfig`
- Delete the associated test files and test cases that tested the
now-removed paths
- ~3800 lines deleted

**Commit 2: Rename Simplify → StmtSimplify**

- Rename `src/tirx/transform/simplify.{h,cc}` → `stmt_simplify.{h,cc}`
- Rename C++ identifiers: `Simplify` → `StmtSimplify`, `SimplifyConfig`
→ `StmtSimplifyConfig`
- Rename FFI keys: `"tirx.Simplify"` → `"tirx.StmtSimplify"`,
`"tirx.transform.Simplify"` → `"tirx.transform.StmtSimplify"`
- Update Python wrappers and all call sites (~40 files)
- Clarifies that this pass operates on statements (distinct from
expression-level `arith::Analyzer::Simplify()`)

## Test plan

- [x] `tests/python/tirx-transform/test_tir_transform_simplify.py` — 52
tests pass
- [x] `tests/python/tirx-transform/test_tir_transform_remove_no_op.py` —
18 pass, 5 xfail
- [x] `tests/python/arith/` — full arith test suite passes
- [x] `tests/python/tirx-transform/` — full suite: 315 passed, 8
xfailed, 1 xpassed (pre-existing vectorize failure unrelated to this
change)
- [x] `pre-commit run --all-files` — all hooks pass
2026-05-26 15:33:40 -04:00
Bohan Hou 859498dc01 [TIRx] Bringup TIRx Infrastructure (#19581)
## 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.
2026-05-18 16:44:43 -07:00
Tianqi Chen 58fc749f27 [REFACTOR] Delete src/support/libinfo.cc; replace with runtime FFI-registry env query (#19477)
## Summary

`support.GetLibInfo` exposed ~30 build-time `TVM_INFO_*` strings (git
hash, LLVM/MLIR versions, every `USE_*` flag). Real callers reduce to
"is this feature enabled?" — better answered at runtime. `USE_CUDA=ON`
does not mean a CUDA device is loadable; runtime discovery is the actual
signal. Git versioning is already tracked via `tvm.__version__`.

Changes:
- Delete `src/support/libinfo.cc`, `cmake/modules/LibInfo.cmake`,
`tests/lint/check_cmake_options.py`, and the `check-cmake-options`
pre-commit hook.
- Delete `tvm.support.libinfo()` (no shim — callers migrate to runtime
discovery).
- Add `tvm.support.detect_active_modules()` which queries the FFI global
function registry for `ffi.Module.create.<kind>` registrations (cuda,
vulkan, opencl). `describe()` now prints active runtimes instead of
CMake build flags.
- Migrate 5 in-tree callers: `_get_targets()` uses `cudnn.exists()` /
`tvm.runtime.enabled()` for CUDNN and Hexagon; `_cmake_flag_enabled()`
is rewritten as a static map from cmake flag names to
`tvm.runtime.enabled()` or FFI-registry probes; `clml_sdk_version()`
uses the existing `relax.get_openclml_version` FFI global;
`test_clml_ops.py` uses the new helper.

After this PR: `src/support/` is header-only.
2026-04-30 07:24:35 -04:00
Shushi Hong b6f67b06db [Docs] Add Python API reference for tvm submodule docs (#19379)
as per title
2026-04-10 14:52:49 -04:00
Akaash Parthasarathy 81889decfa [Fix] Replace str(target.kind) with target.kind.name for Target objects (#18959)
Replace `str(target.kind)` with `target.kind.name` for `Target` objects
since `target.kind` is a `TargetKind` object while `target.kind.name`
yields a string describing the target
2026-04-02 14:17:02 -04:00
Tianqi Chen 141c22fd8a [Refactor] Bring up tirx namespace (#18913)
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
2026-03-19 21:27:54 -07:00
Tianqi Chen 9a8320acbd [LINT][PYTHON] Modernize annotations with ruff UP rules (#18830)
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.
2026-02-27 21:29:47 -05:00
Tianqi Chen d463395706 [REFACTOR][RUNTIME] Phase out legacy contrib runtime backends (#18813)
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
2026-02-23 11:08:03 -05:00
Tianqi Chen 33dcea1686 [REFACTOR][LINT] Modernize ruff config (#18810)
This PR removes the extra lint violations from the codebase so lint
aligns with the latest style
2026-02-23 07:29:21 -05:00
Tianqi Chen aa2e609136 [LINT] Modernize lint to use pre-commit hooks (#18807)
This PR migrates existing lint to use pre-commit hooks
2026-02-22 11:03:21 -05:00
Ruslan Baratov 52e45477de [DOC] Unify CUDA naming (#18797)
Fix CUDA naming in documentation and comments

- Cuda -> CUDA
- cuda -> CUDA
2026-02-19 08:04:00 -05:00
Tianqi Chen 2030db36e4 [REFACTOR][TARGET] Phase out legacy target string in favor of json (#18785)
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.
2026-02-16 16:21:35 -05:00
Tianqi Chen 69018327c9 [chore] Cleanup stale dependencies (#18760)
This PR cleansup stale thirdparty dependencies that are no longer
needed.
2026-02-12 07:15:03 -05:00
Tianqi Chen 0079ff7d1c [REFACTOR][TEST] Replace CompareBeforeAfter for pytest compact (#18711)
This PR refactors test infrastructure by removing the CompareBeforeAfter
base class from tvm.testing and converting all dependent tests to use a
simpler, more explicit pattern.

We need this change as latest pytest do not allow calling fixture as
inner patterns which the previous CompareBeforeAfter depend on.
2026-02-04 11:24:07 -05:00
Siva 8e40211388 [ADRENO][TEXTURE] Texture based lowering (#18523)
Introduces the below features over texture annotation

- Lowering, codegen and runtime for texture.
- image2d_array_t support - Added depth dimension allows more
allocations using texture instead of falling back to buffer when the
texture limits exceeds.
- A comprehensive set of schedules for Adreno textures.
- Texture packing of arbitrary types up to 128 bit (FP16-NCHW8c,
INT8-NCHW16c ...etc.).
- A clBufferDescriptor debug dump controlled by cmake options.
- Pipeline definition for adreno target.


While covering these features the below interfaces or passes or enhanced
which need a review.

- alloc_tensor: VDevice information is passed across these API's. The
way of texture allocation is ```alloc_storage``` allocates buffer/image
objects as requested followed by alloc_tensor being a view of any scope.
This takes care of optimum utilization backing memory across different
image objects or scopes.
- Constants Saving: Handled by adding memory scope section in
executable. This introduces a new header magic to retain the backward
compatibility.
- Static Memory Planing: Mostly port from Relay static memory planner
with mixed mode allocator.

---------

Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
Co-authored-by: Sanjay <sanjs@qti.qualcomm.com>
2026-01-09 14:15:08 -05:00
Guan-Ming (Wesley) Chiu acda952b31 [Relax][PyTorch] Unify tests using shared tvm.testing.assert_allclose (#18522)
## Why

We have the shared assert_allclose func in tests and to use it in every
tests could help persist consistency
2025-11-29 00:51:52 -05:00
Tianqi Chen 3c36ce2ec6 [FFI][REFACTOR][ABI] Rename NDArray to Tensor (#18275)
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.
2025-09-06 14:33:59 -07:00
Tianqi Chen a7a0168be5 [FFI][REFACTOR] Establish tvm_ffi python module (#18226)
* [FFI][REFACTOR] Establish tvm_ffi as a standalone python module

This PR establishes tvm_ffi as a standalone python module.
The ffi is structured as a minimal pip module that can be
directly install by path or url.

examples/get_started provided a minimal example.
This is a major change as we are decoupling tvm_ffi as a
separate package, users need to install tvm_ffi separately.

Thanks to its minimal dependency, tvm_ffi can be easily installed
even just from the source by pip install ./ffi

This change would enable future improvement for library plugins
to have lightweight dependencies by just working on top of
the tvm_ffi, while the main compiler toolchain and runtime
can be layered on top.

* [FFI] Improve traceback setups

This PR improves traceback related setups
2025-08-24 15:46:20 -07:00
Tianqi Chen 4289efa0d5 [REFACTOR][PYTHON] Phase out tvm._ffi and Limited API support (#18020)
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.
2025-05-28 16:52:36 -04:00
Siyuan Feng cd92392d34 [Refactor] Clean up Relay references in the codebase (#17733)
Removing relay references and statements in the codebase.
2025-03-12 09:20:11 -04:00
Siyuan Feng be8e43814a [Refactor] Migrate build API to tvm.compile (#17718)
* tvm.build -> tvm.compile

* relax.build -> tvm.compile

* update
2025-03-09 07:23:52 -04:00
Balint Cristian bbd77ace31 Pick up vector length from 'zvlXXXb' (RVV) mattr for riscv (#17641) 2025-02-19 19:09:06 -05:00
Tianqi Chen c81fccaa2f [REFACTOR] Phase out te.Schedule c++ components (#17662)
* cleanup schedule c++

* remove vitis ai

* remove VERILATOR

* remove aocl and sdaccel

* remove opengl

* remove microdev and antlr

* remove frontends

* fix

* Cleanup relay related legacy components

* fix

---------

Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
2025-02-17 17:03:09 -05:00
Tianqi Chen 9f846bda5b [REFACTOR] Phase out te.schedule python components (#17658)
* [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>
2025-02-16 17:13:10 +08:00
Tianqi Chen ccaa534b2c [REFACTOR] Phase out relay python components (#17656)
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>
2025-02-15 13:48:28 -05:00
Siyuan Feng f717c5655c [Refactor] Phase out microTVM (#17554) 2024-12-10 08:43:05 -05:00
Mengshiun Yu 2a87c4cfc0 [BYOC][NNAPI] Add NNAPI backend for BYOC (#17385)
* [BYOC][NNAPI] This PR intorduce NNAPI to TVM

This PR introduces a new BYOC backend for Android Neural Networks API (NNAPI),
enabling execution of neural networks on custom accelerators. This feature adds
a new codegen and runtime for NNAPI, supporting operations such as element-wise
ops, nn.dense, and nn.conv2d for CNN model with static shape.

Co-authored-by: Ming-Long Huang <mlhuang@pllab.cs.nthu.edu.tw>
Co-authored-by: HMZ <mzhuang@pllab.cs.nthu.edu.tw>
2024-09-23 21:42:37 -04:00
Ruihang Lai 8db545dddd [ROCm] hipBLAS integration (#17290)
This commit integrates hipBLAS into TVM. The minimum ROCm version
requirement is 6.0.

Co-authored-by: Lesheng Jin <leshenj15@gmail.com>
2024-08-22 13:33:04 -04:00
Masahiro Hiramori 7c9969bbdf Remove and replace deprecated distutils.util.strtobool() (#17185)
remove and replace deprecated distutils.util.strtobool
2024-07-23 11:33:06 -04:00
Eirene Pandi c4e6f96386 [TOPI] Add dense schedule for fp16 and fp32 using gemm (#17091)
Add a new schedule for the dense operator
based on the gemm algorithm.
2024-07-09 11:57:18 +03:00
Andrei Hutu cab54e0dee [SME][TOPI] Add conv2d NHWC SME fp32 schedule (#17003)
This commit adds a scalable `arm_cpu` conv2d NHWC schedule for fp32 which generates SME instructions by using the tensor intrinsics introduced in #16921.

Alongside the SME schedule, the logic of the TE schedule `schedule_conv2d_gemm_native()` for both non-scalable and scalable vector implementations has also been translated into the new TIR schedule. This means that the TE compute definition `compute_conv2d_NHWC_hybrid()` is now compatible with both the original TE schedules (e.g. `schedule_conv2d_NHWC_hybrid()`) and the newly introduced TIR schedule `schedule_conv2d_NHWC_hybrid_TIR()`. The corresponding TOPI test has been extended to reflect that.
2024-05-28 17:30:21 +01:00
Andrei Hutu ac9a943c4d [TOPI][Testing] Enable conv2d NHWC fp16 topi testing for arm_cpu (#17007)
This commit adds fp16 test cases to the conv2d NHWC TOPI schedules for `arm_cpu`.
Following the example of #8529, the numpy reference conv2d output is computed in fp32 instead of fp16, while the absolute tolerance varies for each test case according to the size of the summed axis and the output's largest element.
2024-05-22 11:01:02 +01:00
Luke Hutton b49468ddf1 [SME] Introduce scalable fp32 dense schedule (#16921)
This commit adds a new scalable fp32 dense schedule that calls SME intrinsics according to the SME RFC: https://github.com/apache/tvm-rfcs/pull/107.

Currently the schedule does not make use of predication, meaning the output from the matmul compute must be copied in a subsequent compute stage. This will be removed once support for predication is added.
2024-05-15 11:28:16 +01:00
Eric Lunderberg 1d4b9ea5c3 [UnitTest] Use pytest's scope='session' for tvm.testing.parameter (#16930)
Prior to this commit, the `tvm.testing.parameter` utility defined a
fixture with the default `scope="function"`.  However, this prevents
use of these parameters as arguments for other fixtures that are
themselves cached using pytest.  Since these are parameters, not large
values that would be expensive to compute, there is no downside to
caching them at the pytest level.

This commit updates the scope of fixtures generated using
`tvm.testing.parameter` to use `scope="session"` instead of the
default `scope="function"`.
2024-05-14 04:27:08 +09:00
Steven S. Lyubomirsky 6c701fe5b8 [Unity][Parser] Check well-formedness in the parser (#16569)
* Check well-formedness in the parser

* Correct packed funcs in NN frontend

* Support the check_well_formed optional argument to I.ir_module

* Also check well-formedness in TIR

* Enable normalization for individual Relax functions and PrimFuncs

* Use the error raised by the TIR well-formed checker for the message

* Fix tvmscript test failures

* Whitespace

* Fix errors in verify_well_formed test

* Include a more helpful error message

* Fix TIR test failures

* Address well-formed failures in test_tir_specialize

* Correct well-formedness error in test_tir_analysis_oob

* Correct further well-formedness failures

* Remove __tvm_meta__ from test case to avoid parsing error

* Avoid circular import in entryy.py

* Formatting fixes

* lint fix

* Add pylint exceptions

* Fix whitespace

* Fix more failed test cases

* Catch inappropriate use of decl_function instead of segfaulting

* Fix test_lower.py

* Mark purity in test_relax_2d_buffer_allocation.py

* Mark purity in test_dma_builtin.py

* Remove __tvm_meta___ from test_tir_usmp_analysis_extract_bufferinfo.py

* Suppress well-formed check in test_tir_transform_convert_blocks_to_opaque.py

* Remove __tvm_meta__ in test_tir_usmp_algo.py

* Remove __tvm_meta__ from more USMP tests

* Fix incorrect var in test_tir_transform_storage_flatten.py

* Remove all remaining instances of __tvm_meta__

* Fix purity error in test_dataflow_pattern.py

* Fix purity error in test_ast_printer

* Fix test_arith_domain_touched example

* Okay to set check_well_formed to True in test_tir_analysis_identify_mcmcpy

* Define variable in test_tir_analysis_oob

* Typo fix

* Add explanatory comment to test case

* Define the undefined vars in test_tir_transform_common_subexpr_elim

* Exception no longer necessary in test_tir_transform_inject_rolling_buffer

* Remove unnecessary check exemption in test_tir_transform_convert_ssa

* Avoid checking exemption in test_inject_ptx_ldg32

* Note special case in test_distributed_transform_propagate_sharding

* Exempt well-formed error in dlight/test_benchmark

* Exempt well-formedness errors in test_ethosu/, mostly uninitialized vars

* Whitespace

* Include non-CUDA GPUs in IsScheduledOnGPU

* Fix thread binding bug by changing thread binding var dtype

* Include overrides in test_runtime_builtin_paged_attention_kv_cache.py

* add exemptions in test_ethosu/test_replace_conv2d

* Add more ethosu exemptions

* More exemptions for ethosu tests

* Remove unused reference

* Indicate purity in test_transform_rewrite_cuda_graph

* Indicate purity in test_transform_normalize

* Reorder MergeSharedMemoryAllocations in GPU codegen

* Add target parameter for FP8StorageLegalize and FP8ComputeLegalize

* Don't re-import Target in tvm/tir/transform/transform.py
2024-03-21 15:37:18 -04:00
Luke Hutton af0c038f2e [SVE] Add codegen support for scalable buffer accesses (#16696)
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>
2024-03-14 11:48:21 +00:00
Eric Lunderberg 3ec0ca5b0b [Disco] Expose functions to query the per-worker device/rank (#16639)
In addition to the PackedFunc `"runtime.disco.worker_id"`, which
returns the worker ID wrapped in a `ShapeTuple`, this commit adds
`"runtime.disco.worker_rank"`, which returns the worker ID without
wrapping, and `"runtime.disco.device"`, which returns the device for
each worker.

The unit test added in this commit simulates loading of model weights
through a parameter transformation function.
2024-02-26 19:06:15 +09:00