119 Commits

Author SHA1 Message Date
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 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
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 95d1268982 [REFACTOR] Introduce and modernize FFI system (#17920)
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.
2025-05-06 19:18:33 -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
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 09a0ea0cd4 [FIX][TVMC] Fix the mixed precision conversion pipeline
Fixed the mixed precision conversion pipeline issue.
2025-01-26 10:21:48 -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
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
Krishna Bindumadhavan 5645c52c6d [Marvell BYOC]: Marvell AI Accelerator Integration - Phase 1 (#16570) 2024-02-16 22:46:23 +08:00
Liam Sturge ff36526a4c [CI] Cut pytest-lazy-fixture (#16512)
Pytest-lazy-fixture doesn't work with version pytest==8.0.0

The following error occurs:

   AttributeError: 'CallSpec2' object has no attribute 'funcargs'

This has been raised as an issue on the pytest-lazy-fixture project,
but the repository is not in active development. See

https://github.com/TvoroG/pytest-lazy-fixture/issues/65

This patch removes the use of the library to resolve the problem.
2024-02-05 19:23:09 +00: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
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
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
Aleksei-grovety 6cf5b0928a [microNPU][ETHOSU] Fix compiler attributes types (#15159)
The bool type is replaced with tvm::Bool to pass these parameters to tvmc via the command line.
2023-06-29 09:27:39 +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
masahi 0dae36048e [CI] Update CPU image to install PyTorch (#14847)
* [CI] Update CPU image to install PyTorch

* use link-static to prevent symbol conflict problem

* tlcpackstaging -> tlcpack

* disable tvmc pth tests for now

* fixed skip marker
2023-05-20 18:09:41 +09:00
gaoxinge 0e24aa7c8b [FIX] fix typo in comment (#14787) 2023-05-17 07:21:10 +05:30
Luke Hutton f8f7bc8946 [CI] Fix broken model link (#14458)
The path to the `cnn_s_quantized.tflite` model used in some TVMC tests
was changed in https://github.com/ARM-software/ML-zoo/pull/56. The link
has been updated with the new path.

I noticed the same model was used from a different link in the
codebase - using this link to avoid duplication and fix S3 error.
2023-04-04 10:04:44 +01: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
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
Junru Shao c452e6966c [TVMScript] IR Fragment Printing (#13742)
This PR introduces support for TIR fragment printing.
Fragment printing makes it possible to print TIR fragments in the text
format consistency with TVMScript PrimFunc/IRModule printing.

This PR still preserves the legacy ReprPrinter format by introducing an
API `LegacyTIRPrint` for TIR PrimExpr. This method is used in
AutoScheduler and TIR CSE for full backward compatibility.
2023-01-13 22:36:47 -05:00
Christopher Sidebottom bd37515cfb [Tests] Replace pytest.main with tvm.testing.main (#13717)
This ensures that if you want to run a specific test script then at least it's reasonably consistent and as people copy test files they'll use the new function 😸
2023-01-06 21:54:07 -08: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 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
Matveenko Valery 574794e915 [OpenCL] Enable OpenCL for GPU tests (#12490)
* Add opencl target in test build script

* Fix fp16 test and compile test for opencl

* fix lint

* Fix relay OpenCL texture tests

* Fix lint

* Enable relay OpenCL tests

* Fix opencl relay texture tests

* fix lint

* Remove OpenCL gtest variable

* Fix unbound variable

* Skip tests that are not supported in CI

* fix lint

* Add path for opencl gtest directory

* Fix opencl gtests include directory

* Enable OpenCL googletest. Fix bug in opencl timer test

* testing fix for build cpp tests

* update googletest git version for opencl tests build

* update cmakelist

* Update CMakeList

* Update CMakeList

* Disable opencl googletests

* update Opecnl.cmake

* fix Opecnl.cmake

* Apply comments. Remove xfail decerator for opencl tests. Now specific tests are skipped in the environment script

* minor code changes

* apply comments

* apply comment

* skip test in ci by decorator

* fix pytest skipif warnings

* Fix skipif for opencl gtests
2022-09-10 04:01:53 +09:00
Leandro Nunes 2d36e46007 [CI][AArch64] Mark tests to be skipped due to torch crash (#12730)
Some integration tests are not being run on CI due to the
configuration of the machine with onnx and torch not calling
the integration tests script.

This patch skips two more tests failing with the error message
below:

```
"OSError: /.../torch/lib/libgomp-d22c30c5.so.1:
cannot allocate memory in static TLS block"
```
2022-09-08 14:49:55 +01: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
Leandro Nunes d54c0651ec [Torch][AArch64] Skip test_load_model___wrong_language__to_pytorch (#12660)
This patch makes test_load_model___wrong_language__to_pytorch to be
skipped in AArch64 due to a bug that can be reproduced when enabling
Integration Tests in machines with Torch installed in TVM.

```
The error message seen is:
OSError: /usr/local/lib/python3.7/dist-packages/torch/lib/
libgomp-d22c30c5.so.1: cannot allocate memory in static TLS block
```

While the test needs further investigation, it is being set as
skipped so other tests can be enabled and not to regress and allow
time for the investigation to be made.

This relates to the issue described in #10673.
2022-08-31 12:51:21 -07:00
Mehrdad Hessar c15cc5ef6d [Target] Remove deprecated parameters from target (#12416)
* remove depricated parameters in target

* lint

* fix cpp tests

fix

* remove more configs in test files

* address comments

* fix error

* fix hexagon

* fix micro tutorial

* fix integration tests

* fix hexagon

* lint

* fix unittest

* fix readme

* fix assert executor in target

* address comments

* fix tutorials

* fix hexagon target

* fix tutorial

* fix for tutorials

* hexagon
2022-08-24 13:02:27 +09:00
Leandro Nunes d271678388 [CI][AArch64] Skip libgomp failures in integration tests (#12554)
Some integration tests are failing when running in CI machines that
have torch installed (validated only in AARch64 for now), with an
error message related to libgomp, similar to the one above:

OSError: /.../dist-packages/torch/lib/libgomp-d22c30c5.so.1: cannot
allocate memory in static TLS block

As part of enabling the integration tests in AArch64, I'm marking this
tests as skipped, so that tests can start executing and don't regress
while we take time to investigate these specific failures.
2022-08-23 14:03:04 +01:00
Leandro Nunes 48a8cbd57d [ONNX] Fix test to disable default ONNX frontend constant folding (#12532)
In TVM ONNX frontend, constants are folded by default, which makes `test_load_model__onnx` to fail because it is looking for "params" that were already converted into constants.

This patch fixes the test to disable constant folding so that we can assert that "params" in the model are present as expected.
2022-08-22 16:48:59 +01:00
Christopher Sidebottom 1de32b587b [Target] Add Target Parser for Arm(R) Cortex(R) M-Profile CPUs (#12319)
This implements an initial Target Parser which uses the same logic as
the CMSIS-NN compiler flags to update the features and keys of the `c`
and `llvm` `Target`s.

Refactoring of the CMSIS-NN logic will be in a separate patch.
2022-08-12 16:31:34 +01:00
Luke Hutton 759a648cd5 [TVMC] Only load extra targets when there are workspace pools (#12253)
After #11427, `tvmc compile` wouldn't work for external codegens that
don't have a `Target` registered by `TVM_REGISTER_TARGET_KIND`. Such
external codegens can be expected to have no workspace pools and may not
always have a target associated as their implementation predates this
mechanism. While it is likely a `Target` is specified for all external
codegens in the future, we should still support external codegens
without an associated `Target` until this is enforced.

Co-authored-by: Chris Sidebottom <chris.sidebottom@arm.com>
2022-08-02 10:10:40 +01:00
Dhruv Chauhan 75ec1cffa9 [TVMC] Workspace Pools Parameters (#11427)
* [TVMC] Workspace Pools Parameters

Attributes from tvmc are now passable into the created PoolInfo objects
inside WorkspaceMemoryPools. This is passed in to relay.build that get
attached to IRModule attribute.

* [TVMC] Workspace Pools Parameters

Address comments, fix linting. Testing improved.
Change-Id: Iea79329b6b9ec1cbc51e5c293449bf6dd43b00c5

* [TVMC] Workspace Pools Parameters

Update workspace pools test naming
Change-Id: Ib698d6248be1e6f44340f27db3641c985bc5c5d8

* [TVMC] Workspace Pools Parameters

Add test for parameter overrides.

Change-Id: I67d5470dcfbfbc9ab27f34e20a9269d2070193ca

* [TVMC] Workspace Pools Parameters

Rebasing over #10189
Updates to the way a WorkspaceMemoryPool object is created
Change-Id: I1f0e1d240343af311ddb3ed5c564cc1ab329f463

* [TVMC] Workspace Pools Parameters

Fix linting, fix CI
Change-Id: If75f8709ac4ad925655eca54b3e5c1bb09d025e8

* [TVMC] Workspace Pools Parameters

Add mcpu and mattr to target registry for cmsis-nn
Change-Id: I15257b8d01624c071c738cab6d12ecb84ed6cb16

* [TVMC] Workspace Pools Parameters

Added test for override on single pool when multiple pools are present
Updated functionality of parsing multiple attributes
Change-Id: I2c0745051b7a923dd7f75040bfb89bbc99376a11
2022-07-25 10:43:00 +01:00
Philipp van Kempen 4d95f2c9c5 TVMC: Add new text/relay frontend (#10941)
* TVMC: Add new text/relay frontend

This feature enables passing a textural representation of a relay module to the tvmc command line.

Example: `tvmc compile relay.txt --target c --runtime=crt --executor=aot --executor-aot-unpacked-api=1 --pass-config tir.disable_vectorize=1 -f mlf`

Currently it is not possible to supply parameters as it is mainly intended to be used for testing certain relay functions or operators. In the future (with minor changes to the tvmc frontend api) params could be passed via an additional i.e. `params.bin` file

This commit also adds minimal unit testing of the added feature.

Resolve PR comments

TVMC: add warning if relay frontend is used

* [TVMC] populate parameters with random values instead of ones

* [TVMC] Relay frontend: do not populate input tensor buffers if --input-shapes is provided

This prevents that the constants inputs are used for Constant folding,
thus changing the complexity of the model.

If there would be a way, to distinguish between model inputs and parameter this
workaround would not be required.

* [TVMC] Relay frontend: check provided file contents before calling tvm.parser.fromtext()
2022-07-19 16:21:19 +01:00
Eric Lunderberg 9aaae28bbc [TIR] HoistExpression, generalization of HoistIfThenElse (#11592)
* [TIR][Arith] Use non-inlined bindings when proving conditional

* [TIR][Arith] Recognize Var when used as a literal constraint

* [TIR][Arith] Added simplification of constrained if_then_else op

This feels like it should definitely be part of RewriteSimplify, but
that will require making CanInlineLet be a virtual function.

* [TIR] Implemented HoistExpression transformation

This is a generalized form of HoistIfThenElse, which can also hoist
Let bindings, or portions of conditional expressions.  This will be
used in upcoming changes to separate compute loops into a slow loop
that handles edge cases and a fast branchless loop.

* [TIR] Expressed HoistIfThenElse as special case of HoistExpression

* Lint fixes

* Fixed breakage in tvmc unit test that relied on pass type

* More accurate handling of kUsingBlockVar

Didn't correctly reproduce previous behavior.  In addition to
preventing hoisting of expressions that use a block
variable (e.g. threadIdx.x), should also prevent hoisting of
expressions across a "thread_extent" AttrStmt.

* Updated comment for HoistExpression pass

* Fix linting error
2022-06-24 16:15:17 -07:00
Luke Hutton 8bf6cd5800 [TVMC] Fix tvmc run when using rpc (#11757)
* [TVMC] Fix tvmc run when using rpc

As described in #11707, the RPC mechanism does not support
objects of type Map which breaks the use of tvmc run when using
RPC after #9889. This commit intends to workaround this issue by
providing a fallback to the old implementation when RPC is being
used. Further, a test has been provided to help prevent this
regression in the future.

Change-Id: I70c1863d00098270e27c08ba834a3587e9132d69

* fix lint

Change-Id: I958cf4e19988d047bdd2e02f6475b9f70afe80c8
2022-06-20 13:48:15 +01:00
Jason f942d19788 [TVMC] Fix error while compile paddle model with tvmc (#11730)
The tvmc command will throw a error while the passed path of model is not exist, But for PaddlePaddle model, it contains 2 file model_name.pdmodel and model_name.pdiparams, we only pass the prefix like inference_model/model_name.

This pr is same with https://github.com/apache/tvm/pull/11108 
Since the origin PR didn't update for a long time, I send this new PR
2022-06-15 10:02:04 -07:00
Philipp van Kempen e19cf20054 TVMC: Allow to overwrite TVM_CONFIGS_JSON_DIR via environment variables (#11623)
If a non-default location for the build directory is used, e.g. set via TVM_LIBRARY_PATH
we need to provide the user a way to overwrite CONFIGS_JSON_DIR as well.
2022-06-08 14:21:29 +01:00
Eric Lunderberg b885362c36 [CI] Refactor of tvm.testing.requires_* annotations (#11313)
* [CI] Improved skip messages when using @tvm.testing.requires_*

Previously, the same message was given regardless of why a test
couldn't be run.  This has been split up into separate checks for TVM
cmake options in `config.cmake`, enabled targets in `TVM_TEST_TARGETS`
environment variable, and checks for available hardware.

* Refactor to specify repeated feature marks, compile-only markers

* Fixed lint errors

* Import from contrib, not from a different import

* Removed use of requires_llvm() as a list of marks

* Corrected mark from requires_gpu to requires_cuda

* Adding missing "not"

* Added USE_CMSISNN as a requirement for corstone300.
2022-06-03 14:03:08 -07:00
Florin Blanaru 553eb1acd0 [tests] add utility to replace direct call to pytest.main (#11393) 2022-05-23 11:52:22 -07:00