209 Commits

Author SHA1 Message Date
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 543e64dbb1 [FFI][REFACTOR] Cleanup tvm_ffi python API and types (#18277)
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
2025-09-07 10:38:50 -04: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 e7c04f554b [Refactor] Introduce base Executable class and tvm.compile interface (#17710)
This refactor introduces a base Executable class and a `tvm.compile`
interface that can be used to compile both TIR and Relax programs.

`tvm.compile` will return an Executable object that can be used to call
either TIR or Relax functions.
2025-03-07 08:00:48 -05:00
Bohan Hou fc6775e770 [REFACTOR] move build flow from C++ to Python (#17665)
This PR moves build flow from C++ to python, enables more developer productivity and readabilities
2025-02-20 14:47:06 -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
Siva bf77a5ba94 [RUNTIME][CLML] Profiling options enabled for CLML (BYOC via JSON Runtime)
Graph debug runtime to modifications to acommodate profiling through
BYOC external calls. Updated TVMC interface to add more formats while
profile dump. Added CLML helpers that can rebiuild CPP clml sources
from profile dumps. CLML runtime profiling is now controlled by runtime
profile flag.
2025-02-11 12:59:11 -05:00
Siva 92caa2f7e6 [TVMC] Bug fix
Multiple sub graphs of same BYOC is possible.
Append them to the same dump file and also ignore any errors while dump.
2025-01-26 10:22:09 -05:00
Siva 09a0ea0cd4 [FIX][TVMC] Fix the mixed precision conversion pipeline
Fixed the mixed precision conversion pipeline issue.
2025-01-26 10:21:48 -05:00
Siva ed2c26abae [ADRENO][WINDOWS] Windows build dependencies for Adreno target
This brings in basic support for Adreno target on Windows platforms
and the tools we generally use for Adreno
2025-01-26 09:32:52 -05:00
Siyuan Feng f717c5655c [Refactor] Phase out microTVM (#17554) 2024-12-10 08:43:05 -05:00
MNGanesan 988255e6fa Compiled with Default Target(LLVM) and Built with USE_MRVL=ON (#17455)
* [Frontend][ArgParse] Compile with default(LLVM) target and build with BYOC(#17454)

It is a unique use-case to check the default target(LLVM), though TVM is built with BYOC(MRVL-ON)
The config of Codegen(BYOC) contains default values for configuration/options, it is extracted
during _generate_codegen_args. In command line processing, validate_target_args checks if there
are add-on options and it expects that particular target to be given explicitly in command line.
Here, it is test for default (LLVM) path only, hence validate_target_args need to ignore the codegen's
configuration for default target.

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* [Frontend][ArgParse] Compile with default(LLVM) target and build with BYOC(#17454)

It is a unique use-case to check the default target(LLVM), though TVM is built with BYOC(MRVL-ON)
The config of Codegen(BYOC) contains default values for configuration/options, it is extracted
during _generate_codegen_args. In command line processing, validate_target_args checks if there
are add-on options and it expects that particular target to be given explicitly in command line.
Here, it is test for default (LLVM) path only, hence validate_target_args need to ignore the codegen's
configuration for default target.

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* [Frontend][ArgParse] Compile with default(LLVM) target and build with BYOC(#17454)

It is a unique use-case to check the default target(LLVM), though TVM is built with BYOC(MRVL-ON)
The config of Codegen(BYOC) contains default values for configuration/optons, it is extracted
during _generate_codegen_args. In command line processing, validate_target_args checks if there are
add-on options and it expects that particular target to be given explicitly in command line.
Here, it is test for default (LLVM) path only, hence validate_target_args need to ignore the
codegen's configuration

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* [Frontend][ArgParse] Compile with default(LLVM) target and build with BYOC(#17454)

    It is a unique use-case to check the default target(LLVM), though TVM is built with BYOC(MRVL-ON)
    The config of Codegen(BYOC) contains default values for configuration/optons, it is extracted
    during _generate_codegen_args. In command line processing, validate_target_args checks if there are
    add-on options and it expects that particular target to be given explicitly in command line.
    Here, it is test for default (LLVM) path only, hence validate_target_args need to ignore the
    codegen's configuration

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* [Frontend][ArgParse] Compile with default(LLVM) target and build with BYOC(#17454)

        It is a unique use-case to check the default target(LLVM), though TVM is built with BYOC(MRVL-ON)
        The config of Codegen(BYOC) contains default values for configuration/optons, it is extracted
        during _generate_codegen_args. In command line processing, validate_target_args checks if there are
        add-on options and it expects that particular target to be given explicitly in command line.
        Here, it is test for default (LLVM) path only, hence validate_target_args need to ignore the
        codegen's configuration

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* [Frontend][ArgParse] Compile with default target(LLVM) when built USE_MRVL=ON(#17454)

This is a use-case of invoking TVMC with default target though it is built with MRVL_ON.
In command line processing, validate_target_args checks if there are add-on options
derived from the default arguments of codegen/BYOC and it expects that particular codegen
to be given explicitly in command line. However, certain codegen's can have default target alone,
in that case codegen optios are not extracted there by relaxing the validation

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* [Frontend][ArgParse] Compile with default target(LLVM) when built USE_MRVL=ON(#17454)

This is a use-case of invoking TVMC with default target though it is built with MRVL_ON.
In command line processing, validate_target_args checks if there are add-on options
derived from the default arguments of codegen/BYOC and it expects that particular codegen
to be given explicitly in command line. However, certain codegen's can have default target alone,
in that case codegen optios are not extracted there by relaxing the validation

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* [Frontend][ArgParse] Compile with default target(LLVM) when built USE_MRVL=ON(#17454)

    This is a use-case of invoking TVMC with default target though it is built with MRVL_ON.
    In command line processing, validate_target_args checks if there are add-on options
    derived from the default arguments of codegen/BYOC and it expects that particular codegen
    to be given explicitly in command line. However, certain codegen's can have default target alone,
    in that case codegen optios are not extracted there by relaxing the validation

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

---------

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>
Co-authored-by: M N Ganesan <muthusamynam@marvell.com>
2024-10-25 09:52:57 +01:00
Siyuan Feng 9e865b4b8f [Docs] Introduce Relax API and move legacy part to standalone page (#17286)
* [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
2024-08-22 18:16:56 -04: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
MNGanesan 02fe0c5f0d [Frontend][ArgParse] Pass default values to target compiler(#13264) (#17014)
* [Frontend][ArgParse] Pass default values to target compiler(#13264)

    BYOC Compiler's Config node defines the target compiler's
    command line options, along with default values. This change
    extract the default values from config node, while constructing
    target options for codegen/target compiler.
    Added test case for this feature as well.

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* [Frontend][ArgParse] Pass default values to target compiler(#13264)

    BYOC Compiler's Config node defines the target compiler's
    command line options, along with default values. This change
    extract the default values from config node, while constructing
    target options for codegen/target compiler.
    Added test case for this feature as well.

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

* Lint Fix

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>

---------

Signed-off-by: M N Ganesan <muthusamynam@marvell.com>
Co-authored-by: M N Ganesan <muthusamynam@marvell.com>
2024-06-26 18:28:11 +08:00
Eric Lunderberg b2204ae698 [IR] Default to empty attributes, instead of NULL (#16745)
* [IR] Default to empty attributes, instead of NULL

Prior to this commit, the default `DictAttrs` for an `IRModule`,
`tir::PrimFunc`, `relax::Function`, and `relay::Function` was a null
value.  At each callsite, the absence of a `DictAttrs` needed to be
treated as equivalent to an empty `DictAttrs`.  In C++, this typically
was done using the `foo->GetAttr` helper function, but in Python it
needed to be checked explicitly.  That is, every callsite needed to
check `if func.attrs is not None and attr_name in func.attrs`, rather
than only checking `if attr_name in func.attrs`.

Since most functions would have at least one attribute to specify the
global symbol, these bugs would often surface when working on
unrelated changes.

This commit changes the default attribute dictionary from
`NullValue<DictAttrs>()` to `DictAttrs()`.  This avoids having two
separate representations of an object without any attributes, and
allows the `if attr_name in func.attrs` pattern in the Python API.

* Remove no-longer-needed checks on attrs being present

* Fix up unit tests

* More unit test fixes

* Undo erroneous find/replace

* A few more unit tests

* Provide `DictAttrs.get`
2024-03-25 13:07:56 -05:00
Krishna Bindumadhavan 5645c52c6d [Marvell BYOC]: Marvell AI Accelerator Integration - Phase 1 (#16570) 2024-02-16 22:46:23 +08:00
Eric Lunderberg 7336debd8d [Bugfix][Relax] Remove call to tvm.build for empty TIR module (#16561)
Prior to this commit, if a lowered `IRModule` does not contain any TIR
functions, `tvm.relax.build` provided an empty `tir_mod`, which caused
a segfault during TIR compilation.  This could occur when
`tvm.relax.build` is called without an explicit target argument, for a
module that does not define any virtual devices.

This commit updates the `_filter_tir` utility function to return
`None` if there are no TIR functions, rather than an empty
`IRModule`.  In addition, checks for an empty `IRModule` are added to
`tvm.build` and `TIRToRuntime`, so that a similar failure mode would
raise an exception rather than producing a segfault.
2024-02-14 07:07:27 -06:00
tqchen c40d96b59e Merge remote-tracking branch 'upstream/main' into unity 2024-01-11 12:12:15 -05:00
Eirene Pandi 8e67e2a3a1 [TVMC] Add tvmc flag to print ir before and print ir after named pass (#16261)
Add `--print-ir-before` and `--print-ir-after` options  to print IR before and after a named pass from the compiler's command line driver.
2024-01-11 11:52:25 +00:00
Yong Wu 1c35c39264 [Unity] Add Relax multi-device e2e cases (#15823)
* [Unity] filter out non-GPU primfuncs in default_gpu_schedule

* Add relex heterogeneous e2e case

* Remove get_prim_func_device

* Update test cases

* Fix flake8

* fix lint

* Add test case for change of default_gpu_schedule

* fix comment
2023-12-20 13:52:56 -08:00
Siva cde83e1088 [TVMC] enable dumping imported modules too (#15779)
Now we can dump the imported modules source too like device code.
2023-09-25 16:56:45 +01:00
Krzysztof Parzyszek f9e6018cfe [Runtime] Make export_library parameters after file_name keyword-only (#15658)
This makes the code a bit more readable at a little cost.
2023-09-01 23:45:46 -07:00
Andrei Hutu 9ff74fb13f [TVMC] Add tvmc flag to print compilation time per pass (#15349)
Added a new flag `--print-pass-times` for tvmc compile to provide debugging information for tvmc users using `PassTimingInstrument`. Also added a test to check the printing of timing results.
2023-07-25 16:44:16 +01:00
sergio-grovety fa223b077f [Bugfix][TVMC] Fix tvmc option for printing which operators are offloaded to the Ethos-U (#14994)
Lines from the initial Relay, which don't correspond to the relay.Call now are printed to the output with the --dump-offloads option enabled. Thus we get rid of the incomprehensible gaps in the line numbers of the initial relay, which occurred before.

---------

Co-authored-by: Sergey Smirnov <89378719+sergey-grovety@users.noreply.github.com>
Co-authored-by: Arina.Naumova <naumova@grovety.com>
2023-06-19 10:17:04 +03:00
Yong Wu aa7d2bff6b [CI] Modify test cases to accommodate the CI upgrades (#14651)
* [CI] update all the images

* Update test_config

* Fix llvm_codegen_test err

* Update with newly built images

* update pylintrc

* Update i386 build

* Don't use ninja for i386

* Update torch tests

* Debug i386 platform

* check ec2 instance type for i386

* Remove gluoncv ssd example

* Update pylint

* Update test images

* Skip torch jit trace issue for arm

* Fix pylint

* update tests

* update i386 build

* update s3.py to skip non-existing files

* update pylint

* Update pylint

* Fix tests

* update clang-format to 15

* update tests for clang-format-15

* run with newly images

* skip oom test for i386

* Upgrade for DGL sample

* fix black

* Ignore a warning in doc

* New run with newly images

* Use newly generated tlcpackstaging images
2023-05-05 11:54:41 +01:00
Balint Cristian 515583ce28 [AutoTVM] New rank-binary loss_type for the new xgboost >= 2.0.0 behaviour (#14468)
New rank-binary loss_type for the new xgboost >= 2.0.0 behaviour

Signed-off-by: Balint Cristian <cristian.balint@gmail.com>
2023-04-11 08:23:05 -04:00
Philipp van Kempen ffc1fc0116 [TVMC] Allow selecting a subset of tasks to be used in tvmc tune (#12525)
This adds a `--tasks` flag to the `tvmc tune` command to filter the lists of tasks to be tuned. See examples below.

## Motivation

- As auto-tuning can be quite time consuming, it is often desirable to cut down the number of tuned tasks in a session.
- If the tuning session was canceled halfway through, it would be a bad idea to start from scratch. Instead continue with the last untuned task
- Some tasks have more impact on the model performance than others, thus we should be able to train some tasks longer than others

## Examples

1. Use `--task list` to show which tasks are available for tuning
```
$ tvmc tune toycar.tflite -o out.txt --task list
Available Tasks for tuning:
  0. Task(func_name=dense_nopack.x86, args=(('TENSOR', (1, 640), 'int16'), ('TENSOR', (128, 640), 'int...
  1. Task(func_name=dense_pack.x86, args=(('TENSOR', (1, 640), 'int16'), ('TENSOR', (128, 640), 'int16...
  2. Task(func_name=dense_nopack.x86, args=(('TENSOR', (1, 128), 'int16'), ('TENSOR', (128, 128), 'int...
  3. Task(func_name=dense_pack.x86, args=(('TENSOR', (1, 128), 'int16'), ('TENSOR', (128, 128), 'int16...
  4. Task(func_name=dense_nopack.x86, args=(('TENSOR', (1, 128), 'int16'), ('TENSOR', (8, 128), 'int16...
  5. Task(func_name=dense_pack.x86, args=(('TENSOR', (1, 128), 'int16'), ('TENSOR', (8, 128), 'int16')...
  6. Task(func_name=dense_nopack.x86, args=(('TENSOR', (1, 8), 'int16'), ('TENSOR', (128, 8), 'int16')...
  7. Task(func_name=dense_pack.x86, args=(('TENSOR', (1, 8), 'int16'), ('TENSOR', (128, 8), 'int16'), ...
  8. Task(func_name=dense_nopack.x86, args=(('TENSOR', (1, 128), 'int16'), ('TENSOR', (640, 128), 'int...
  9. Task(func_name=dense_pack.x86, args=(('TENSOR', (1, 128), 'int16'), ('TENSOR', (640, 128), 'int16...
```

2. Filter the list of tasks to be tuned:

```
# Only tune a single task (index 5)
tvmc tune toycar.tflite -o out.txt --tasks 5

# Tunes tasks starting with index 6
tvmc tune toycar.tflite -o out.txt --tasks "6-"

# Tune tasks 1,4,5,6,8,9
tvmc tune toycar.tflite -o out.txt --tasks "1,4-6,8-"
```

## Tests
I added a basic unit test for the `filter_tasks` utility in `tests/python/driver/tvmc/test_autotuner.py`.

## Open Questions
- ~~While the (truncated) string representations of AutoTVM tasks are quite helpful to pick the correct tasks, using AutoScheduler the tasks can not really be distinguished from each other (only by index). Is there a way to get similar information from AutoScheduler tasks?~~
2023-03-30 14:13:08 +01:00
sergio-grovety da8335378a [TVMC][microNPU] tvmc option for printing which operators are offloaded to Ethos-U (#13212)
Added an option to tvmc and Ethos-U for printing to console or to the file which operators from the initial graph are offloaded to Ethos-U and which aren't. It forms line-by-line output of initial model IR, indicating which operations ported to Ethos-U.

Compiler option "--target-ethos-u-dump_npu_functions_coverage" has been replaced by more generic "--dump-offloads" with the same meaning.


## Usage
```
# output to console:
tvmc compile --target=ethos-u,cmsis-nn,c \
    --dump-offloads=- \
    ........

# output to file:
tvmc compile --target=ethos-u,cmsis-nn,c \
    --dump-offloads=<file path> \
    ........
```

## Example output:


...
Total number of operators and distribution by targets
Total: 211
target1: 198
target2: 10
generic: 3

'target1        <-     target2.qnn_conv2d'
'target1        <-          %0 = qnn.conv2d(%tfl.quantize, %v_param_1, ...'
'target1        <-          %1 = nn.bias_add(%0, %v_param_2, axis=3);'
'target1        <-          %2 = qnn.requantize(%1, meta[relay.Constant]...'
'target2        <-     target2.reshape'
'target2        <-          %3 = reshape(%2, newshape=[1, 1001]);'
'generic        <-     %4 = nn.pad(%3, -128f, pad_width=[[0, 0], [1, 1]...'
...
2023-03-27 20:26:10 +01:00
Philipp van Kempen f4520c4f15 [TVMC] Improve --desired-layouts functionality (#14272)
this aims to make the `--desired-layout` argument more powerful based on the previously merged changes from #14010 by introducing two new features:

1. Allow passing multiple arguments to `--desired-layout` instead of only one, to specify one layout per transformed operator specified in `--desired-layout-ops`. (Number of arguments has to bei either 1 or match the number of transformed operators)
2. Optionally, you can now specify a non-default kernel layout as follows: `NHWC:HWIO`

Example Usage: `tvmc compile … --desired-layout nn.max_pool2d qnn.conv2d --desired-layout-ops NCHW NHWC:HWIO`

I also added unit tests for the new use-cases.

### Known Limitations:
* It would make sense to specify individual kernel layouts for regular convolutions and depthwise ones. However since both are usually implemented as generalized `nn.conv2d`, we can not transform them individually. Are there any good workarounds for this?
* The arguments of `--desired-layouts` have previously been checked for validity during cmdline parsing (e.g. only NCHW and NHWC are allowed) which is not possible anymore. Should I add a regular expression for that?
2023-03-16 10:30:59 +00:00
Elen Kalda f6b75792c5 [TVMC] Fix logging in TVMC (#14175)
Three logger related changes in this patch:
* Currently we don't set the output stream on the Python logger, so
it defaults to sys.stderr, which means we only get some logger output
when the command fails. So set the output stream to sys.stdout
* Currently we can add -v flag to anywhere in the command line for
tvmc compile, but only between tvmc and run/tune for run and tune.
Unify the behaviour such that we can add the flag anywhere on the
command line.
* Set the effective upper bound of -vs to 3 as 4 could result in
NOTSET which would not output anything.
2023-03-15 15:31:31 +00:00
Siva 52292cfa60 [TVMC][TRANSFORMS] ToMixedPrecision transform support with custom options enabled (#14010)
Adds new command line options:
* `--mixed-precision` - Enable mixed precision conversion
* `--mixed-precision-ops` - List of operators to be converted to mixed precision
* `--mixed-precision-calculation-type` - Calculation precision type
* `--mixed-precision-acc-type` - Accumulator precision type
    
Additionally: 
* `--desired-layout-ops` - The list of operators to be transformed with desired layout.
2023-03-09 09:53:45 +00:00
Luke Hutton a15ade30b1 [TVMC] Add option to dump TIR code to file (#14186)
* [TVMC] Add option to dump TIR code to file

Dump TIR code after the final phase of lowering (phase
3 from https://github.com/apache/tvm/blob/665dd413bc85d14f7836324daf7cc0dd9281c85a/gallery/how_to/extend_tvm/low_level_custom_pass.py#L152)
before codegen. This is done by running a pass to capture the TIR
module as it is not saved during the build. The result is saved to
file similar to the other code dumps.
2023-03-06 09:39:07 +00:00
Mehrdad Hessar 05cbe329d8 [microTVM]Enable TVMC micro with AoT Executor (#14077)
This PR enables AoT Executor for tvmc micro compilation.
2023-03-02 08:34:20 -08:00
Elen Kalda 62a69a6c93 [TVMC] Stop printing a wall of warnings with tvmc tune (#13882)
A simple tvmc tune command currently results in a huge wall of warnings
about target_host parameter being deprecated, even when the user hasn't
provided a target-host cmd line argument.

We can prevent that happening from just not providing the default
target-host to tvmc. Also, ensure that when the user does provide
target-host, we print the warning once, not 500 times.
2023-02-03 10:32:54 +00:00
MNGanesan 26d3244fb8 Enhance the --help message of composite target (#13842)
Presently --help for vitis displays the target and option string,
it has no description.  Eg:  target vitis-ai dpu<class 'str'>

This can be made more meaningful by fetching the description from
the config node of the target.  Eg: Vitis AI DPU identifier

Signed-off-by: MNGanesan <mnganesan@yahoo.co.uk>
2023-01-25 13:49:55 +00:00
Siva 1f40b925a5 [TOOL][NATIVE] Android native application for deploy and run (#13791)
* [TOOL][NATIVE] Android native appliction for deploy and run

This application helps as a reference for verifying and integration of
TVM compiled models on Android targets natively independent of RPC setup.

tvmc will be used to for compiling tuning and to run it before deployment.

This PR also covers
 * Enabling clml for tvmc compilation tool.
 * Graph runtime api "get_output_info" to return output tensor specification
   similar to "get_input_into"
 * This tool adds and enabled 3rdparty dependency "cnpy" to deal with npz files.

* Update apps/cpp_rtvm/README.md

Co-authored-by: Egor Churaev <egor.churaev@gmail.com>

* Update apps/cpp_rtvm/README.md

Co-authored-by: Egor Churaev <egor.churaev@gmail.com>

* * review comments.

* * proof reading

* Update apps/cpp_rtvm/README.md

Co-authored-by: Egor Churaev <egor.churaev@gmail.com>

* * review

Co-authored-by: Egor Churaev <egor.churaev@gmail.com>
2023-01-24 16:11:54 +03:00
Luke Hutton 23ade0c14b [TVMC] Global pass context for compile and tune (#13309)
* [TVMC] Global pass context for compile and tune

Comes as a followup from conversations in #13216. By making the pass
context a global value for both `compile` and `tune` commands, we can
ensure the pass context is exactly as the user expected and also
test components such as `convert_graph_layout` under a pass context
suitable for testing (e.g. add instruments). With this change, it
becomes the users responsibility to ensure the PassContext they
select is suitable for the passes that will be run. By default,
`opt_level` remains as 3 so current workflows that do not alter the pass
context from the command line / TVMC Python API should not be affected.

Change-Id: I7a601daf6fbe664f77bce1b45efeb7ca29f621b3

* fix vitis-ai test and typo

Change-Id: I04f5bd031ae4717825f42e373bcb0e1e2c1c9d90
2022-11-10 13:49:20 +00:00
Luke Hutton 79093a1718 [ETHOSN] Consolidate target string usage (#13159)
* [ETHOSN] Consolidate target string usage

Removes support for a deprecated target string. The deprecation warning
has been around for a couple of releases now so it should be safe to
remove. The target to use moving forward is: `ethos-n -variant=n78 ...`

Refactored direct use of a driver stack target string in the testing
infrastructure to use the same string we expect users to provide. This
simplified some of the code in codegen and hopefully avoids confusion
in the future.
2022-11-08 10:19:52 +00:00
Luke Hutton 4ecf303695 [TVMC] Apply constant folding when converting layout (#13216)
This commit ensures that constant folding is applied when a desired
layout is selected during compilation. It ensures that
`layout_transform` operations are removed where possible so that
pattern matching for BYOC backends can work effectively.

A test has been added to check this regression.
2022-11-02 14:15:00 +00:00
Yineng Zhang 209e77c18b [tvmc] add instruments for PassContext (#13136) 2022-10-20 12:14:48 +01:00
Christopher Sidebottom c9002509f6 [Target] Print deprecation warning before canonicalisation in build module (#12747)
Hopefully fixes #12742, as the warning should only be printed when a user passes `target_host`, in the current case if the user passes `None` as `target_host` it'll be processed by `canon_target_map_and_host` which seems to always produce a `target_host` and thus triggering the warning despite the user doing nothing wrong.
2022-09-15 14:03:21 -07:00
Luke Hutton da48e13b66 [TVMC] Run module once by default (#12713)
* [TVMC] Run module once by default

Currently executing `tvmc run module.tar` will run the input model
twice. For benchmaking this is to be expected as the first run is used
to prime caches etc before taking a measurement. However, this seems a
bit unintuitive to have as default, especially when benchmarking is not
always intended. In this sense, this commit aims to amend the
number of runs for the default: `tvmc run module.tar` to a single run.

After inspection, this seems to be down to the use of the `.benchmark()`
method which runs (1 + repeat * number) executions in total. This means
that at least two runs are required (i.e. when repeat=1, number=1). It
also seems that it is only necessary to benchmark the model when
`--print-time` has been set from the CLI POV. From the python interface
POV, benchmarking is always run, but this may not always be necessary.

This commit makes use of the `.run()` method to singularly execute the
model by default. From the CLI this will be used when `--print-time` is
set to False whereas from the python interface this will be used when
`benchmark=False`. Otherwise, the `.benchmark()` method will be used
as before. Complementary to this change `repeat`, `number` and
`end_to_end` parameters are only used when either `--print-time` or
`benchmark` are set to True - and the documentation has been updated to
indicate this.

Change-Id: I18a38a9d430d660264f7fce5caf0779aa059fed3

* improve documentation with number of exectuions when benchmarking

Change-Id: Iecf557594420fcc9f3abcec5ce7d952db2c94271
2022-09-06 21:10:36 +01:00
Yulv-git bdcfa01eae [Fix] Fix some typos (#11503)
Fix some typos in src/.

Co-authored-by: driazati <driazati@users.noreply.github.com>
2022-08-19 16:31:30 -07:00