44 Commits

Author SHA1 Message Date
Tianqi Chen 302aaf9f96 [IR] Rename Var name_hint field to name (#20016)
Rename the reflected local `Var` field from `name_hint` to `name` and
update its typed C++ consumers. Preserve distinct named-node APIs and
the Python constructor keyword compatibility path, while making `.name`
the sole stored Var property. Upgrade legacy compact JSON records for
current and pre-unification Var schemas.

Validation: full runtime/compiler build, focused C++ Var copy-helper
test, focused Python IR/Relax/TIRx/script tests, Vulkan codegen syntax
build, touched-file pre-commit checks, and `git diff --check`.
2026-07-17 05:34:31 +08:00
Tianqi Chen c717c5b217 [IR][Relax][TIRx] Unify Var identity (#20004) 2026-07-15 17:12:40 +08:00
Tianqi Chen adf8d6a463 [TIRx] Phase out duplicate Var type_annotation (#19944)
## Rationale

TIRx variables use inherited `ExprNode::ty` as their single semantic
type. Retaining a primitive handle surrogate erases the distinction
between scalar values, typed pointers, and true opaque pointers, then
forces later passes and code generators to reconstruct information that
the IR already owns.

## Changes

- Remove the duplicate reflected `Var::type_annotation` state and
preserve exact `PrimType` or `PointerType` through construction,
visitors, transforms, specialization, builders, printers, and code
generation.
- Keep scalar-only boundaries explicit through `PrimExpr`, `PrimVar`,
and `PrimType`; pointer-capable values remain general `Expr` or `Var`.
- Keep helper boundaries no broader than their contracts: TE tensor
variable indices use `PrimVar`, while expression deep equality recurses
through general `Expr` only where pointer-bearing `Call` arguments
require it and does not generalize private arithmetic subclasses.
- Keep core statement reflection typed as `Expr`, name general
reinterpret targets as `target_ty`, and preserve exact pointer calls in
the general vectorization path with explicit scalarization behavior.
- Delete `PrimType::Handle()` and `PrimType::IsHandle()`. True opaque
pointers use `PointerType::VoidPointerTy()`; TVMScript renders the
canonical global type as `T.handle`, standalone values as `T.handle()`,
and scoped void pointers with a keyword-only storage scope.
- Make `CodeGenSourceBase::SSAGetID` a single `Type` boundary across
source backends, without a separate primitive-type or runtime-dtype
variant.
- Keep WebGPU semantic argument classification type-aware: storage
buffers are identified from `PointerType`, POD arguments from
`PrimType`, and only the final `FunctionInfo` launch ABI is serialized
to `DLDataType`.
- Preserve exact pointer semantics at runtime boundaries, including
access pointers, packed calls and returns, external calls, storage
rewrites, and target-specific lowering.

## Migration guide

- **Variable types:** In C++, replace `var->type_annotation` with
`var->ty`; in Python, replace `var.type_annotation` with `var.ty`. The
result is the exact `Type`: scalar variables carry `PrimType`, while
pointer variables carry `PointerType`.
- **Scalar boundaries:** Use `PrimVar` and `PrimExpr` for variables and
expressions that are semantically scalar. When starting from a general
view, narrow explicitly with `var.as_or_throw<PrimVar>()` or
`expr.as_or_throw<PrimExpr>()`. Keep pointer-capable fields and call
arguments as `Var` or `Expr`. A default-constructed `PrimVar` is
nullable, so construct local scalar variables explicitly, for example
`PrimVar i("i")`.
- **Opaque pointers:** Replace `PrimType::Handle()` with
`PointerType::VoidPointerTy()`. Replace `IsHandle()` tests with explicit
`PointerType` inspection; use `PointerType(element_type, storage_scope)`
when the pointee type is known instead of erasing it to a runtime handle
dtype.
- **TVMScript handles:** Use `arg: T.handle` for a global void-pointer
annotation and `arg = T.handle()` for a standalone value. Use
`T.handle(storage_scope="shared")` for a scoped void pointer. Typed
pointers use forms such as `T.handle("float32")`, `T.handle("float32",
"global")`, or `T.handle("float32", "shared")`. Legacy
`T.handle("void")` input remains parse-compatible, but the printer
canonicalizes it to `T.handle` (or the keyword-only scoped form).
- The separate `tirx.type_annotation` intrinsic used by access-pointer
APIs is unchanged; this migration removes only the duplicate variable
field.

## Validation

- Complete native C++ test executable: 122/122 passed, including
`IRF.CountVar`.
- Relax binding-rewrite suite: 12/12 passed, including transferred-user
bookkeeping.
- Canonical typed/void/scoped TVMScript handle printer and round-trip
checks: 5/5 passed.
2026-07-04 21:54:04 -04: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
Tianqi Chen 120812e9ac [REFACTOR][Relax] Phase out PrimValue and Relax expression wrappers (#19891)
This PR lets Relax expressions directly take `PrimExpr` values without
requiring the explicit `PrimValue` wrapper, continuing the Relax IR
unification work by removing Relax-specific leaf/base expression layers.

Summary:
- Remove `LeafExpr` / `LeafExprNode` and use direct expression-node
checks where needed.
- Converge Relax expression typing onto the shared IR `Expr` base.
- Remove the `PrimValue` node wrapper while keeping `relax.prim_value` /
`R.prim_value` as conversion helpers that return existing `PrimExpr`
values unchanged.
- Register direct `PrimExpr` handling through exact concrete node
dispatch, aligned with the `tirx` expression visitor list and excluding
arith iter-map intermediate nodes.
- Inline the private Python primitive conversion helper into public
`relax.prim_value`.
- Handle direct `PrimExpr` values in frontend scalar paths without
assuming a `.value` field on non-immediate expressions.
2026-06-26 07:18:04 -04:00
Tianqi Chen 1e1920bcbd [REFACTOR][IR] Unify PrimExpr type mechanism to PrimType instead of DataType (#19875)
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
2026-06-24 21:31:47 -04:00
Tianqi Chen 1bb5cf6102 [REFACTOR][IR] Unify StructInfo and Type (#19853)
## 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
2026-06-21 10:12:12 -04:00
Tianqi Chen 1240649257 [FFI][REFACTOR] Direct structural APIs to tvm-ffi (#19661)
## 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
2026-06-03 18:57:05 -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 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
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 17113f8216 [REFACTOR] Formalize namespace for all objects (#18101)
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
2025-07-01 07:19:23 -04:00
Tianqi Chen 10e66934ec [REFACTOR] Phase out the RelaxExpr.checked_type in favor of struct_info (#18078) 2025-06-18 16:39:06 -04: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
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 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
Eric Lunderberg 02f48828e4 [FFI] Re-introduce the boxed primitive values (#17257)
* 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
2024-08-12 08:36:17 -04:00
Tianqi Chen 11be832620 Revert "[FFI][RUNTIME] Introduce runtime boxed types for int/float/bool" (#17252)
Revert "[FFI][RUNTIME] Introduce runtime boxed types for int/float/bool (#16183)"

This reverts commit 5f22be4d83.
2024-08-07 12:19:13 -04:00
Eric Lunderberg 5f22be4d83 [FFI][RUNTIME] Introduce runtime boxed types for int/float/bool (#16183)
* [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
2024-08-05 09:19:20 -04:00
Ruihang Lai 2dbab710e8 Merge branch 'main' into 'unity' 2023-11-16 10:30:20 -05:00
Siyuan Feng 3f3473e2d5 [TIR] Enhance Python Type Annotations for TIR Expr (#16083)
This PR enhances the Python annotations for the TIR expr,
adding class member variables annotations.
2023-11-08 10:54:59 +08:00
Siyuan Feng 3f1347cbd4 [Unity] Enhance Python Annotations for Relax Expr (#16075)
This PR enhances the Python annotations for the Relax Expr, adding
class member variables annotations and improving docstring
2023-11-06 00:15:44 -08:00
tqchen 90c64c6dce [MERGE] Merge main into unity 2023-09-06
NOTE: use the original webgpu impl to make sure webgpu is stable.
2023-09-06 10:45:29 -04:00
Anirudh Sundar Subramaniam 04ee895d8d [IR] Use structural equal for Range equality (#15664)
This PR adds a small change to verify equality of `tvm.ir.Range` as a
structural equal. This assumes that in most cases, comparing two
`Range`s means to compare its `min` and `extent` as opposed to the
actual Range object handle.
2023-09-05 12:29:11 -05:00
tqchen 9eafe7c022 [MERGE] Merge main into unity 2023-06-13
Merge remote-tracking branch 'upstream/main' into unity-staging
2023-06-13 21:23:33 -04:00
Krzysztof Parzyszek 00126b0484 [IR,TE,TIR] Use f-strings for string formatting, NFC (#14990)
* [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
2023-06-01 19:35:41 -05:00
tqchen 9a3bbc1d37 [MERGE] merge main into unity 2023-05-20 2023-05-20 17:29:56 -04:00
Eric Lunderberg f4a7eaebd6 [TIR][TVMScript] Added format/parsing of subroutine calls (#14889)
* [TVMScript] Cherry-pick module.other_func syntax from unity

* [TIR][TVMScript] Added format/parsing of subroutine calls

Similar to `module.relax_func(args)` syntax used when parsing Relax
functions, this allows `module.tir_func(args)` to be used when parsing
TIR PrimFuncs.
2023-05-20 09:13:37 -05:00
Siyuan Feng 540ba28f5c [Unity] Relax TVMScript Parser. (#13932)
This PR adds the TVMScript parser/ir_builder support based on the blockbuilder.

Co-authored-by: Ruihang Lai <ruihangl@cs.cmu.edu>
Co-authored-by: Junru Shao <junrushao1994@gmail.com>
Co-authored-by: Tianqi Chen <tianqi.tchen@gmail.com>
Co-authored-by: Yuchen Jin <yuchenj@cs.washington.edu>
Co-authored-by: Steven S. Lyubomirsky <slyubomirsky@gmail.com>
Co-authored-by: Yong Wu <yongcale@gmail.com>
2023-04-01 15:31:36 -04:00
Yuchen Jin e68ef58c84 [Unity] Basic StructInfo Analysis and Expr construction (#13916)
[Unity] Basic StructInfo Analysis and Expr construction.

This PR adds struct info analysis and expr support.
These are logics to construct the IR node and perform
struct info related analysis.

Testcases are added to cover the IR node construction
and related struct info analysis checks.

Co-authored-by: Tianqi Chen <tianqi.tchen@gmail.com>
Co-authored-by: Altan Haan <altanh@cs.washington.edu>
Co-authored-by: Andrew Liu <andrewlliu@gmail.com>
Co-authored-by: Hongyi Jin <3231950289@qq.com>
Co-authored-by: Jiawei Liu <jaway.liu@gmail.com>
Co-authored-by: Junru Shao <junrushao1994@gmail.com>
Co-authored-by: Lesheng Jin <34279105+LeshengJin@users.noreply.github.com>
Co-authored-by: masahi <masahi129@gmail.com>
Co-authored-by: Prakalp Srivastava <prakalp@octoml.ai>
Co-authored-by: Ruihang Lai <ruihangl@cs.cmu.edu>
Co-authored-by: Siyuan Feng <Hzfengsy@sjtu.edu.cn>
Co-authored-by: Steven S. <Lyubomirsky slyubomirsky@octoml.ai>
Co-authored-by: Sunghyun Park <49998730+sunggg@users.noreply.github.com>
Co-authored-by: Yixin Dong <ubospica@gmail.com>
Co-authored-by: Yong Wu <yongcale@gmail.com>
Co-authored-by: Ziheng Jiang <ziheng@apache.org>
2023-04-01 15:31:36 -04:00
Junru Shao e77a6d1a05 [TVMScript] Introduce PrinterConfig (#13831)
This PR introduces `PrinterConfig`, a systematic way to configure
TVMScript printer without having to set global flags.

This PR enables more customization of printer behavior. More
specifically, now any TVM’s object in python, as long as it
inherits from `Scriptable`, it automatically gains two methods:
- `.script(tir_prefix=...)`
- `.show(...)`
2023-01-24 06:54:29 -08:00
Junru Shao da99e9d1b5 [TVMScript] Use TVMScript for all TIR Printing (#13795) 2023-01-18 08:42:24 -05:00
Wuwei Lin 44c35dcd96 [TVMScript] Fix parsing int64 loop with optional loop start (#13068)
* [TVMScript] Fix parsing int64 loop with optional loop start

* Update scope_handler.py

* fix range

* fix
2022-10-14 14:59:57 -07:00
Mark Shields fb99383c43 [Relay] Re-run PlanDevices after LowerTE to flow new memory scope constraints. (#9613)
* [Relay] Re-run PlanDevices after LowerTE to flow new memory scope constraints.

This PR:
 1) Makes PlanDevices consider lowered calls when solving device domain constraints.
 2) Connects the storage scopes on PrimFunc parameters (encoded in their Buffer data
    Var type annotation PointerTypes storage_scope fields) to the memory_scope
    fields of the SEScopes which PlanDevices unifies over.
 3) Allows new device_copies to be inserted on the arguments and results of lowered
    calls so as to acount for any memory scope mismatches which are now apparent.

[device_planner.cc has main changes, rest is secondary.]

In the short term we'd like to use this machinery to flow memory scope choices made
during lowering back out into the overall Relay program. In the longer term we'd
also like to be able to use memory scopes to influence the lowering of
yet-to-be-lowered functions (or lowered functions which have yet to been scheduled,
a distinction now possible with TensorIR).

 - Memory scope constraints can flow both out of and in to PrimFuncs
   introduced by LowerTE. In TIR memory scopes are represented by
   'storage scopes' on the PointerType type annotations on TIR Buffer data
   variables.
    - It is straightforward to extract memory scopes from PrimFuncs by
      looking at the PrimFunc's buffer_map. We do this is 'phase 1' of
      PlanDevices, which collects all the device constraints implied by
    - However, pushing memory constraints in to PrimFuncs is more challenging
      due to buffer aliasing. This aspect is still experimental.

 - Allow device_copies to be inserted for both arguments and
   results of PrimFunc calls, on the assumption PlanDevices has
   already established a consistent device assignment prior to
   lowering and any new mismatch is required to match up memory scopes.
   We use the new 'free' on_device annotations to implement this.

Coming along for the ride:

 - To make unit tests of mixed Relay/TIR functions possible needed
   to be able to supply a checked_type to GlobalVar since that's currently
   the only way to give a Relay type to PrimFuncs.

 - Use GenSym to get unique var names in ANF & partial eval so easier
   to diff debug output between passes and connect program fragments
   back into the overall program. Relying on pretty-printing to
   automagically unique-ify var names is certainly cute but until we
   have better span support is very hard to work with.

 - Realized both dead_code.cc and fold_constant.cc would
   happily move values into a different lexical virtual
   device context since device_planner.cc was being
   'clever' and eliding on_devices for let-bound values
   when there's no change. Fixed so that every let-bound
   value has an on_device. Will be much better after
   https://github.com/apache/tvm-rfcs/pull/45 is implemented.

 - Make build -Werror clean for clang-12 (mostly move fixups).

 - Address post-submit comments from #9693.

* [checkpoint] thread safe GenSym
2021-12-14 11:16:26 -07:00
Tristan Konolige fdfc7eb887 [TVMSCRIPT] Attach span information to tir nodes in tvmscript (#6910) 2020-12-04 20:03:08 -05:00
Jared Roesch f13fed55cf [Format] Convert all Python code w/o CI (#6448)
* Add black setup

* Tweak pyproject.toml

* Fix syntax issues

* Fix

* Tweak

* Black all Python code
2020-09-11 22:17:24 +09:00
Tianqi Chen 0465ffda26 [REFACTOR][TIR][API-Change] Range/IntSet API style consistency. (#5953)
- Range::make_by_min_extent -> Range::FromMinExtent
- Update the APIs in IntSet to use CamelCase
2020-06-28 16:02:06 -07:00
Tianqi Chen 6027412bcb [IR] Update the type_keys to reflect the code-org (#5074) 2020-03-15 15:39:47 -07:00
Tianqi Chen e03164159c [TIR] Introduce tir::PrimFunc (#5070)
This PR introduces tir::PrimFunc which will be used as the TIR function
container in the unified IR.

Also streamlined the function attributes a bit further.
- All common attributes are under tvm::attr
- TIR specific attributes are under tvm::tir::attr and comes with a tir prefix
- Use stl_style for attributes for now
2020-03-14 14:26:23 -07:00
Tianqi Chen ec86d7f1a2 [REFACTOR] Streamline Function Attr interface. (#5045)
* [REFACTOR] Streamline Function Attr interface.

There has been quite a few recent changes that depends heavily on
the function attr interface. This PR streamlines that interface by introducing
two APIs that covers most of the usages.

- GetAttr which gets a typed object for a given key
  - HasNonzeroAttr is a quick helper that calls GetAttr to quickly check an attribute
- WithAttr that creates a new function object with the given attr
  - The API comes with copy on write optimization to avoid multiple copies
  - We deliberately pick the prefix With(instead of Set) to indicate this
    function does not mutate the original input.

On the python side:
- We allow read access via func.attrs (which is a DictAttr)
- func.with_attrs to create a new instance with updated attrs.

We also get rid of the small wrapper functions and make sure the API centered around
the GetAttr and HasNonzeroAttr interface.

This PR also changes the function construction to follow the new convention.

* Address review comments

* Address review comments

* Fix doxygen path
2020-03-12 08:33:47 -07:00
Tianqi Chen 08338dd5f8 [REFACTOR][PY] Establish tvm.te and tvm.driver (#4900)
- Move the related files to tvm.te
- Move build_module.py to tvm.driver
2020-02-17 18:47:09 -08:00
tqchen b787ffa34a [REFACTOR][PY] Establish tvm.tir
- Move related files into the corresponding location as in C++
- Keep the top-level TVM API backward compatible to make minimum changes in topi
2020-02-13 20:24:01 -08:00
Tianqi Chen a566161147 [REFACTOR][PY][API-CHANGE] establish tvm.ir, migrate corresponding files (#4862)
* [REFACTOR][PY][API-CHANGE] establish tvm.ir, migrate corresponding relay files.

This PR establishes tvm.ir and migrates the corresponding relay
files into the new folder.

API Change:
- relay.Module -> tvm.IRModule

* Update with ADT

* Migrate transform

* address comments

* Migrate module

* Migrate json_compact

* Migrate attrs

* Move LoweredFunc to stmt temporarily

* temp migrate container

* Finish migrate container
2020-02-11 20:01:36 -08:00