## 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.
## 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.
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
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.
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 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.
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.
* [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
[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 formalizes the namespace for all object registered so
we do not have object that sits on root namespace
Also fixes the Visitor style in TensorMapNode
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 modernizes the FFI foundation of the project and introduce
a new minimal and lightweight module [tvm ffi](https://github.com/apache/tvm/tree/refactor-s3/ffi)
based on our lessons in the past few years. It implements a modern
version of the [Unified Packed and Object RFC](https://github.com/apache/tvm-rfcs/blob/main/rfcs/0097-unify-packed-and-object.md)
that unifies the packed function call and object systems.
Summary of the change:
- A dedicated clean Any/AnyView that can store strong and weak
references of items
- Function(previously PackedFunc) system built on top of the Any/AnyView
- A minimal C API that backs the overall calls. We are stabilizing the
API with a goal to bring clean, stable FFI conventions for both compiled
and registered code
- A rewrite of core python binding and generated code based on the module
- Update existing code and test cases to the new module
- Latest dlpack support
The new module brings many benefits thanks to the cleaner design,
to name a few:
- Any can support both POD types(int) and object types.
- Containers (e.g. Array) can now also contain Any value, e.g. now
`Array<int>` is supported, no need for boxed types
- Error handling now upgrades to object-based, allowing cleaner
traceback across languages
- Map now preserves insertion orders
- Path toward isolated stabilize minimum core ABI/API foundation module
- Type traits based design that cleanly defines how values interact
with Any system
- Automatic conversion of different types based on traits if needed
Because FFI upgrade is at heart of the project, the change touches every
component of the system. Importantly, this is an upgrade of the ABI so the
change is not backward compatible. The code compiled under the old
FFI won't work under the new one. We did provide example ABI translation
(e.g. LegacyTVMArgValueToFFIAny) functions for compatibility.
The PR tries to leave files in their old places while creating redirections.
The goal is to have the first milestone landed and infrastructure in place,
so we can do further refactors to complete features and cleanup legacy code
as trackable PRs. As of now, python binding and compiled code are under the
new convention while RPC and some other bindings still relies on legacy ABI
translation. We will work on upgrades in the coming PRs, including areas such
as reflection, phasing out legacy redirections etc.
* Refactor decorate function to use functools.wraps
Replace decorator package with functools.wraps for simpler and more standard function wrapping. This change removes an external dependency and uses Python's built-in functools module for function decoration.
* Remove decorator package dependency
Remove the decorator package from various installation scripts and requirements files. This follows the previous refactoring of decorators to use functools.wraps, eliminating an external dependency and simplifying the project's package requirements.
* Modify decorate function to support more flexible function wrapping
Update the decorate function to create a wrapper that allows more flexible function decoration. The new implementation passes the original function as the first argument to the wrapped function, enabling more dynamic decoration behavior while maintaining the functools.wraps functionality.
* Remove decorate function and replace with functools.wraps
Remove the custom decorate function from base.py and update multiple files to use functools.wraps directly. This change eliminates the need for a custom decorator implementation and simplifies the codebase by leveraging Python's built-in functools module.
* Remove generic_func.py as upstream did
* Fix dominant issues of decorator
* fix for pickle memoize
* [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>
* [Docs] Introduce Relax API and move legacy part to standalone page
As the TVM project evolves, the Unity strategy has been the recommended
way to use Apache TVM applications. Hence, we are pushing documentation
for the Relax API to the forefront and moving the legacy part to a
standalone page, which may be removed in the future.
* update for ci
* update for ci
* 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
This commit adds support for splitting via the compile-time unknown
constant `vscale`. Two main changes are introduced; they are described
below.
The split scheduling primitive has a new parameter disable_predication
that allows the user to avoid introducing a block-level predicate when
splitting with a factor of `vscale`. This feature is useful when schedule
writers know that the loop they're splitting is a factor of the scalable
vector length for their target. Otherwise, a predicate must be introduced
due to the nature of `vscale`.
CanProve has been extended to prove expressions that use multiple
instances of `vscale`. Known possible scalar values of the `vscale`
intrinsic are iterated over and substituted into the expression. If
the expression holds true for each possible value, we can conclude the
expression true. Currently only support for an SVE target has been
added, but it is possible to extend to other targets as/when needed. If
the analyzer becomes more powerful in the future and is able to deal
with multiple instances of a symbolic value in an expression, this
feature can be removed.
---------
Co-authored-by: Elen Kalda <elen.kalda@arm.com>
Co-authored-by: Neil Hickey <neil.hickey@arm.com>
Removed instances of accidentally repeated words from comments. There
are cases where duplicated words appear legitimately, those cases remain
unmodified.
- Allow `CreatePrimFunc` args to be a mixture of `te::Tensor` and `tir::Var`.
- Integrate `CreatePrimFunc` and `CreateRelaxPrimFunc` into one function.
```python
idx = te.var("idx", dtype="int64")
m = te.var("m", dtype="int64")
n = te.var("n", dtype="int64")
tensor = te.placeholder((m, n), name="tensor")
slice0 = te.compute((idx, n), lambda i, j: tensor[i, j], name="slice")
# use idx as an arg
te.create_prim_func([tensor, idx, slice])
```
* [IR,TE,TIR] Use f-strings for string formatting, NFC
Replace uses of % and .format() with f-strings.
Reformat modified files.
* Rearrange pylint directives for better formatting
Previously, type-checking of a callable arguments, such as to
`tvm.ir.transform.module_pass`, was done using
`isinstance(arg, (types.FunctionType, types.LambdaType))`. This check
can give false negatives for valid python types, such as a bound
method or an instance of a class that implements `__call__`.
This commit replaces the checks with the builtin function `callable()`,
which handles any Python object that can be called using function-like
syntax.
* Added data type pass unification pass to by default promote data types of all indices and shapes to int64 when creating prim func.
* Added some fixes for lowering passes to make it compatible with int64 data type.
This fixes#13330, which was blocking my work to write TIR schedules for microTVM.
I originally thought I'd have to change the function signature of `DomainTouchedAccessMap`, but I couldn't think of a way to do that cleanly. Instead, I changed `extern_primfunc` to use `primfunc.params` to create the buffer lists in the right order.
#13330 should have been caught by `test_tir_te_extern_primfunc.py`, but one of that test's helper functions had the same bug as `extern_primfunc`. I've thus modified `test_tir_te_extern_primfunc.py` to instantiate the input tensors a different way, allowing it to catch regressions of this issue.
* make elem_offset of the buffers created by te.extern a variable
Co-authored-by: Eric Lunderberg <elunderberg@octoml.ai>
* add test
* fix te extern create_prim_func test
Co-authored-by: Eric Lunderberg <elunderberg@octoml.ai>
* [QNN] Disable QNN canonicalization pass.
This commit enables work of TVM without QNN canonicalization pass.
It adds new TOPI ops for QNN + simple compute/schedules.
* added dependence of the qnn::transform::Legalize pass launch on target.
* Added new dense topi operator for the pattern qnn.dense+bias+requantize
* Added support of axis attribute for QNN TOPI ops
* Fixed TOPI compute implementation for qnn.add
* Fixed issue with non zero padding value for qnn.conv2d
* Fixed Bias.add for qnn.conv2d
* Added support of depthwise qnn.conv2d topi operator
* Added support of 1D quantization params in qnn.dequantize
* Added support of qnn.concatenate
* Fixed out of range array access
* Added meta_schedule_original_shape attribute in QDenseAttr and
QConv2DAttr
* Added support of qnn.batch_matmul as a standalone op.
* Added per channel zp in qnn.dense and qnn.conv2d.
* Fixed corner cases like dense+bias+bias+rq.
* Added unit test.
* Removed rq_out_dtype and axis attributes declaration in QConv2DAttra and
QDenseAttrs.
* Changed target x86->Hexagon to disable QNN passes.
* Fixed issue with QDenseAttrs and QConv2dAttrs.
* Fixed build for Cortex-M.
* Removed QDenseAttrs and QConv2dAttrs
* Fix tests after rebase
* Address code review comments.
* [QNN] Add option to disabe QNN passes.
QNN passes are enabled by default. To disable use
disabled_pass=["qnn.Legalize"] in pass config.
* Revert changes of GetPassPrefix interface.
Tianqi Chen