[BYOC] JSON Runtime with DNNL End-to-End Flow (#5919)

* json runtime

* json dnnl WIP

* fix ArrayNode usages

* Support composite functions

* DNNL json runtime: conv2d/add/relu/dense/bn

* add a more complex example

* fix bias memory issue

* rebase to upstream

* merge to metadata module, remove the unused driver

* handle constant

* support composite functions

* support DNNL constant

* clean up

* Simplify dnnl user code

* GetDataSize

* fix dense bug

* improve cmake

* zero copy

* add unit test

* move json to contrib/json

* fix cmake

* lint

* max_digits10 for fp serialization

* only keep base getfunction

* fix lint

* zero copy for all data entries

* address comments

* enable ci

* address comment; fix bug

* address comment

Co-authored-by: Zhi Chen <chzhi@amazon.com>

[CI] Update ci-cpu to the latest (#6031)

[DOCKER] Pin keras version (#6032)

[TARGET] each option of target str should only contain one '=' (#5988)

src/target/target_id.cc ParseAttrsFromRawString L222:
if ((pos = FindUniqueSubstr(s, "=")) != -1)
require option contains only one '='

Signed-off-by: windclarion <windclarion@gmail.com>

fix typos in comments and relay tutorial (#5999)

* [TypoFix]fix typos in comments and relay tutorial

* retrigger

[RUNTIME] if a param not in input, we still consume it's data (#5990)

so the read pointer of stream can move forward

Signed-off-by: windclarion <windclarion@gmail.com>

[PYTORCH]Gather op support added (#6013)

* [PYTORCH]Gather op support added

* retrigger

Remove deplicate line (#6017)

[Frontend][Relay] Add Parser 2.0 (#5932)

Option to specify alternate directory to output build to (#6016)

This is useful when you would like to manage 2 separate builds in the same tvm tree. You can specify a build directory when using make by adding OUTDIR=alternate-build-dir.

Change-Id: I3efed1135343f3903007115ce5dd683ef7bd9e8c

[TEST][FLAKY] test_arith_solve_linear_inequality.py::test_multi_equal (#6014)

[Frontend][MXNet] MXNet frontend support for AMP cast op (#5976)

* amp_cast

* fix test

* more tests

* test more ctxs

* fix doc

* fix typo

* address CR comment

* fix lint

* revert doc change

* Revert "revert doc change"

This reverts commit a410dd5569730ac81af67ddb333c3afbe97eddd7.

* fix doc

* Update relay_pass_infra.rst

Co-authored-by: Ubuntu <ubuntu@ip-172-31-42-138.ec2.internal>

[VTA] Move compiler related registry items to vta/build_module.py (#6012)

Cache object refs in loop partitioner instead of object pointers (#6004)

* Cache object refs in loop partitioner instead of object pointers

Loop partitioner modifies the IR, which can cause TIR objects to
become dead and be destroyed. To avoid working on junk data cache
object references instead of object pointers.

* Fix format/lint errors

[LLVM] Auto-convert shuffle with single index to "extract element" (#6006)

* [LLVM] Auto-convert shuffle with single index to "extract element"

Data types with a single lane are treated as scalars in TVM. On the
other hand, in LLVM there is a difference between a scalar type and
a vector type with a single lane. Because of that, a shuffle with
a single index is equivalent to extracting an element in TIR, but
not in the generated LLVM IR. This patch changes the LLVM codegen
for shuffle to auto-convert single-lane vectors to scalars.

* Try another build

Fix what looks like bizzare copy-paste issue (#6010)

[DOCKER] Only pass pythonpath for ci images (#6005)

Dynamic Tile Op (#5983)

* first working dynamic tile passes first test

* add dyn tile to dynamic_to_static

* fix cpplintt

* respond to review comments. Thanks @siju-samuel

* make dynamic tile compatible with numpy API

Undefuned names: import os for line 324 & import re for line 308 (#6003)

Update main.yml (#6002)

Fix tune_relay_cuda.py (#6001)

[Arith] Inequalities solver (#5618)

[Relay][Frontend][Onnx] Small bug fix for Conv1D imports. (#5995)

* Fix autopad bug in onnx importer for conv1d.

* Fix output shape in test.

* Undo commented out lines oops.

[Target] Use TargetNode::attrs for Target serialization (#5993)

[TFLite] QNN support for TFLite 2.1.0 quantized models (#5848)

* [TFLite] TFLite 2.x parser quantization support.

* Address comments. Fix a bug for depthwise conv

* Added tests for relu, conv, quantize. Address comments.

* Using web-data. Minor refactoring.

* Removing TF hub package

* Trigger CI.

* Handle TFLite input layer naming.

* Addressing reviews.

* Retrigger CI.

[LLVM] VectorType::get with two parameters is deprecated in LLVM 11+ (#5984)

In LLVM 11+ the distinction between fixed and scalable vector types
has become more explicit. Before the introduction of scalable vector
types VectorType::get(e,n) created what is now a fixed vector type.
With the addition of scalable types, it is recommended to use
FixedVectorType and ScalableVectorType classes directly. Alternatively,
there is a VectorType::get that accepts a 3rd parameter indicating
whether the type should be fixed or scalable.
Using the older VectorType::get that implicitly assumes the fixed type
is deprecated and LLVM now generates a warning.

Change calls to VectorType::get to FixedVectorType::get to avoid
compilation warnings.

[Tutorial] Demo showing how to run a pruned 🤗 model. (#5975)

[Target] Migrate data structure of TargetNode (#5960)

[LLVM] Remove redundant function CreateBufferVecPtr (#5982)

The functions CreateBufferPtr and CreateBufferVecPtr do the exact
same thing, so there is no need for both of them to exist. The
latter is only used in place, which further suggests that the
distinction is unnecessary.

fix tvm relay testing tf.py typo error (#5977)

[TOPI] Fix x86 conv2d template when tuning with unpacked layout (#5938)

* fix x86 conv2d and conv2d_transpose template

* address comments

[Relay/TOPI][OP] Add meshgrid op in Relay, TOPI, Pytorch frontend (#5961)

* Add meshgrid op with pytorch importer

* Fix c++ lint

* Fix pylint

* Meshgrid: add scalar test for pytorch, add topi python wrapper

* Add indexing mode attr.

* Add MeshgridAttrs python binding

* c++ lint

[RELAY] Add resnet-3d & Update network definitions for NHWC layout (#5945)

[DYNAMIC] Add Dynamic reshape to a dynamic namespace and add DynamicToStatic Pass (#5826)

* Dynamic reshape passing tests

* Add Dynamic to Static Pass

* rename test file to prevent pytest conflicts

* fix clang build

* add nested dynamic shape test

* remove cuda tests until VM supports dynamic shapes

* rename namespace from dynamic to dyn

* fix lint

* fix lint again

* Remove incorrect doc strings

* remove dynamic behavior from standard reshape

* fix some tests

* merge dynamic and static interfaces in python

* fix missing import

* missed a reference to relay.dyn.reshape

* fix vta example

* respond to review comments

Add MXnNet parser for box_decode (#5967)

Improve docker/bash.sh to handle git worktrees (#5970)

* improve error code when git ls-files fails

* fix docker/bash to handle git worktrees

Print right number of parentheses for LoadNode (#5965)

Stop printing the unnecessary ')' after each LoadNode that didn't
have a matching '('.

Raise an exception when extern function does not return Stmt (#5964)

The function for tvm.te.extern should return either PrimExpr or Stmt,
however there is no check if it actually does so. If it does not, the
result may be a segmentation fault later on. Catch this case early on,
so an informative message can be shown.

Fix small typo in nn.conv2d_gemm_weight_transform (#5925)

* Fix small typo in nn.conv2d_gemm_weight_transform

Change-Id: I7844d898ebf82592f78f478982262ef95f83cc3e

* Add TOPI conv2d_gemm unit tests

Change-Id: I9ed82a68acffcf0dd9720781f8be4aada9d8e6e4

Make first order gradient graphs more efficient (#5959)

Previously, nodes are visited as often as they are used and each time a
derivative is computed. Only at the leaves were the contributions of
everything added. This patch changes this to add at any node that is
used several times.

Fix the meaning of conv{1,2}d_transpose output_padding parameter. (#5758)

* Add output_padding to generic

* Add output_padding to the reference impl

* Add output_padding to arm_cpu

* Add output_padding to the test

* Add output_padding for cuda

* Add output_padding for x86

* Make use of the new output_padding argument in Relay

* Adjust conv2d_transpose Relay test

* Fix lint errors

* Fix the VTA declaration of conv2d_transpose

* support for output padding in conv2d transpose

* some output padding will break IR pass

* Fix new conv2d_transpose test

* Update tophub

* Fix conv1d output_padding too.

* Fix the conv1d_transpose reference function.

* Fix the cuda impl

* fix the topi test for conv1d

* format

* Add tests for conv1d_transpose output_padding and some check that the values are valid.

* Add check in the implementations

* Add checks to the implementations of conv2d

* Make use of the output_padding argument from topi in relay.

* Fix relay tests asking for invalid output_padding

* Fix line length

* Fix vta tests

* Update tophub references

* Trigger CI

Co-authored-by: Thierry Moreau <tmoreau@octoml.ai>

Amendments for gradients (#5941)

* Amendments for gradients

- We fix the dtype handling of consts in generated gradients.
- We add a collapse_sum_to instruction mirroring the collapse_sum_like.
  While for general definitions (potentially dynamic shapes),
  collapse_sum_like is the first choice, when moving to static,
  using collapse_sum_to will greatly simplify the graph.
  (This simplification is not part of the PR.)

* Fix Broadcast rel description in comment

Thank you, @MarisaKirisame

[RELAY][GRAD] handle Tuple/TupleGetItem in first order gradient (#5946)

* handle Tuple/TupleGetItem in first order gradient

* Unify MultiOnes/MultiZeros.

Fix some typo errors in license header (#5956)

Signed-off-by: leonwanghui <wanghui71leon@gmail.com>

[OpenCL] Fix OpenCL get_valid_counts errors due to intrinsic atomic_add (#5857)

* [OpenCL] Fix atomic add used by get_valid_counts

* Rename l -> load, add flag to enable atomics

* Opencl doesn't do data rearrangement

[TIR][ANALYSIS] Refine side effect analysis. (#5954)

[Relay] symbolic max_output_size  (#5844)

* symbolic max_output_size

* pylint

* fix ci

[BUGFIX] Add cuda 11 to contrib.nvcc.find_libdevice_path() (#5902)

[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

[RELAY][VM] Add shape_of instruction (#5855)

add rm xla attributes in tf docs (#5950)

raise right error in tensorflow split op (#5951)

[TIR] Improve Let/LetStmt support. (#5949)

Let/LetStmt are useful primitives to create variable bindings.
While let binding are harmful for simplification and integer analysis,
they are useful for other cases:

- C0: LetStmt is useful to represent a step that has side effect(e.g. call a PRNG)
- C1: Let expression can be used to create deep nested expression for complicated functions.

This PR improves the let support in the following ways:
- Enable vectorization support for let
- Change let simplification strategy to simplify the most trivial case
  while ignore more complicated cases(to avoid deep nest explosion)
- Enhance arith module to handle const bound and modular set for let.

The overall recommendation is to only use Let in the cases when necessary(C0, C1).

[Doc] minor fix for release doc (#5948)

fix string argument mismatch in GraphRuntimeCodegen (#5933)

[TIR][PASS] Remove legacy HoistIfThenElse (#5944)

This pass has not been migrated to the new transform API,
and contains potential bugs per https://github.com/apache/incubator-tvm/issues/5559.
Given that it is not being actively used, this PR remove this pass
from the collection.

Followup PRs are more than welcomed to land a better version that
conforms with the new transform API.

Update date in the NOTICE (#5942)

[TIR][OP][API-CHANGE] Remove CallNode.call_type in favor of attribute. (#5937)

This is a followup refactor for tir::Call.
Now that we have switched call->name to call->op, the function effect property
can be registered through the op itself, so we no longer need the call_type in the CallNode.

- Introduce CallEffectKind to provide a more fine grained categorization of calls.
- Introduce call_pure_extern and call_llvm_pure_intrin to
  allow us to indicate pure calls in those cases.
- Migrate existing usecases to the new API.

add dnnl (#5936)

[CODEGEN][CONTRIB] Various update for CoreML codegen (#5934)

* [CODEGEN][CONTRIB] Various update for CoreML codegen

* fix lint error

[Runtime] Only initialize required module (#5926)

* init required modules

* trigger ci

* trigger ci

Update code_review.rst (#5923)

editorial pass with corrections

Add TupleGetItem to CSE (#5931)

* Add TupleGetItem to CSE

* rename a local variable

[Arith][GPU]Rewrite simplify fix for Vectorized Cooperative Fetching (#5924)

[PatternLang] Don't rewrite expressions used outside of the pattern (#5930)

* Don't rewrite expressions used outside of the pattern

* add comments

[TE] Add LegalizeInvalidAttach to legalize the compute_at location after split or fuse (#5917)

* Add LegalizeInvalidAttach

* lint & typo

* lint & typo

* address comment

* fix lint

refine error (#5929)

[BACKPORT-0.6][Bugfix][Arith] keep div_mode during floordiv simplify (#5922)

Two small fixes to AMDCPU codegen for LLVM 10+ and ROCm 3.5+ (#5920)

- For LLVM 10+ we need to avoid calling Align with 0, or else
  we get a crash.
- For ROCm 3.5+ we need to use code object 3 (the default in LLVM 9+)
  but for ROCm < 3.5 we want the code object 2.
- As we want to separate codegen from the API, we need to add
  a device api query for the version.
  But every one else wants now one, too. (But I only filled it
  in for CUDA for now.)
- I'm throwing in an addition of kMaxRegistersPerBlock for ROCm.
  This was introduced for CUDA in #5898.

Update install.rst (#5858)

* Update install.rst

minor cleanups/corrections

* Update install.rst

Fixed broken link

[Relay][Vm] Some performance improvement to VM (#5901)

* make alignment constant

* tweak copyto and loadscalarint

* some safety check

* x

* lint

* fix

CUDA device API & VerifyGPUCode pass update (#5898)

* Add kMaxRegistersPerBlock device api for cuda

* Add vectorize check to verify_gpu_code

* Lint fix

* Cast fix

[Thread Backend]Fix CPU Thread Binding for Multiple Sockets (#5918)

* Fix CPU Thread Binding for Multiple Sockets

* Backward compatibility

Add MicroTVM tutorial using the STM32F746 discovery board (#5655)

* Add MicroTVM tutorial using the STM32F746 discovery board
with a sample tflite model

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Fix: add a reference to the new turtorials/micro directory

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* fix: Cosmetic, align Micro TVM text with divider

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Fixes to remove warnings, spaces for readability, code blocks

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* remove use of dload in favor of requests for obtaining the TFLite model

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* add setup for CMSIS_ST_PATH
comment out portion of tutorial that will not run without a physical board available

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Fix warning due to ** in python but part of a comment block
The block is commented out since it can only run on device

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Numerous reworks to address feedback.

Within docs/conf.py place the microTVM tutorial prior to the VTA tutorials

Within the micro_tflite
  - rework section headers
  - reorder code so model prep code is all in one place as well as code
    for running on device
  - address indentation feedback
  - remove '' '' usage which I mistakenly thought was getting around a
    sphinx issue involving **

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Change disable_vectorize to use current approach with tvm.transform.PassContext
Change to pull example model from github with download_testdata
Add 2.5K tflite model
Couple of small changes following https://sphinx-gallery.github.io/stable/syntax.html

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* remove use of relay.build_config in favor of PassContext

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Couple of minor 4 space fix ups

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Change to use tvm.transform.PassContext for disable_victorize and disabling FuseOps

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Remove binary module from repo
Change download_testdata back to pull model from linaro server

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Couple of small cosmetic changes. (spaces and extra lines)

Signed-off-by: Tom Gall <tom.gall@linaro.org>
* Convert link to tf docs to examine a tf lite model to use RST syntax

Signed-off-by: Tom Gall <tom.gall@linaro.org>

[TIR][REFACTOR] Deprecate FreeStmt (#5890)

Currently FreeStmt is not being used.
While it can be useful to have an early free hint
we can always use an intrinsic instead of a first class statement.

[PatternLang] Support any index matching for TupleGetItem (#5909)

* support any index matching

* update doc

Fix serialization of inf float value (#5912)

Don't multiply by constant 1 uselessly in dense (#5911)

PyTorch frontend: fix handling of duplicate use of a model weight (#5897)

This happens e.g. in shared input/output embeddings in BERT
or siamese networks.

Thank you @siju-samuel for reporting.

Allow implicit conversion in TVM FFI to tvm::Bool (#5907)

Add Binary Intrinsic ops to TIR Ops in C++ (#5900)

* Add Binary Intrinsic ops to TIR Ops in C++

* clang-format

add a few gradients (#5899)

Rust Refactor Stage 4: Rewrite Rust graph runtime to use new APIs (#5830)

* Port graph-runtime to new API

* --amend

* Fix file lint

* Remove old travis file

* Add @kazum's patch

* Update rust/tvm-sys/src/datatype.rs

Co-authored-by: Andrew <amcharg@gmail.com>
Co-authored-by: Andrew <amcharg@gmail.com>

Fix the python intrin rule (#5895)

[Relay]Allow every runtime module to handle constants (#5885)

* update source module

* address comment

remove fatal (#5888)

[RFC] Improve quantized convolution performance for armv8 architectures (#5754)

* Improve quantized conv2d performance for armv8

Signed-off-by: Giuseppe Rossini <giuseppe.rossini@arm.com>
Change-Id: I3a3d29f5332dd9b3354e8e0dfb24677a521f9c8f

* Add ASF header to conv2d_gemm.py

Change-Id: I33853279e39c849ae1b555a9c91d7557985a0a35

* Run clang-format-10 on c++ files

Change-Id: Ieee22f032e595dabfc1616ab33466fcbf8d94365

* Fix pylint errors/warnings

Change-Id: I435d4d7bca7500db99547f4401fdc0d0995a1ff4

* Fix pylint errors/warnings in topi

Change-Id: I2fc1ad8453e9020072ab967c849df5390c2967b5

* Fix legalizations tests for aarch64

Change-Id: I0a67a49a7849f52ef7d57b9292ce9125bbb7cb2c

* Reintroduce conv2d_nhwc_spatial_pack.arm_cpu and int16 cast

Change-Id: I91b67fabd475e90a9b75f2dd5ecfee851265e0bb

* Switch type of legalization depending on the strategy used

Change-Id: I9a03040a8c40a6cd2658ed14c3751e05a8e19f2b

* Revert last commit

Change-Id: Ice34101e358e3ce8ebfb12c58f73e910ba5de8e8

* Fix the auto-tuner by registering the correct schedules

Change-Id: Id9273688b2620e1ea849ab01b4c46af8fbf37fd0

* Address review comments

Change-Id: Ia1755a0af7b6d159072d9f0c93c932c481101e48

* Improve usability and readability of conv2d_gemm_weight_transform

Change-Id: I3333186bbc2fe4054b58ce15d910e3be7b315482

* Change variable name to weight in Conv2DGemmWeightTransformRel

Change-Id: Ifb5f1f33af7512fe67c6b049b20a42a0bb2d26c9

* Fix clang-10 linting errors

Change-Id: I25ccc844d9cee23766096e1daddb6180abc413a6

* Trigger tests

Change-Id: Id37706fb7cf77a87a3cc817ecf8046297d9ca95a

[TIR][REFACTOR][API-CHANGE] Change Call.name to Call.op(RelayExpr) (#5863)

* [TIR][REFACTOR][API-CHANGE] Change Call.name(string) to Call.op(tvm::Op/RelayExpr)

This PR brings a major refactor to the tir::Call structure.
The current Call structure uses a string field(name) to identify the
function/intrinsic being called. This approach is limited as we start
to expand TIR to be more structured. In particular, we are interested in
the following aspects:

- Type a function and perform better compile time type checking so that we
  can find errors early.
- Register additional properties about an operator, such as:
  - Whether an intrinsic can be vectorized
  - What is the adjoint function of the intrinsic(for tensor expression AD)
  - Whether the operator has side effect.
- Perform specific codegen about an intrinsic if necessary.
- Call into another function in the same module.

The refactor changes the Call.name field to Call.op.
The Call.op field has a RelayExpr type, and we can pass:

- A tvm::Op which represents the corresponding intrinsic.
- A tvm::GlobalVar for calling into another function in the IRModule.

All the current intrinsics are migrated by registering an tvm::Op.
Because the unified IR shares a single Op registry. We use the "tir"
namespace for tir related intrinsics, for example bitwise and is now registered
under `tir.bitwise_and`.

To simplify upgrade, we introduce a `tir.call_extern` intrinsic
that allows us to call into arbitary external function without type checking.
However, we should move towards more type checked variants in the system.

Under the new op design. We should no longer try to pattern match all the
specific intrincis. Instead, we should rely on attr of each Op to do transformation.
For example, the vectorization pass depends on the TVectorizable property of the op,
which can be registered independently.

In this way, we can still grow the number of intrinsics when necessary
without having to change all the passes.

The same rule applies for tensor expression AD. Currently we are performing
AD by pattern match on operators like exp, sin, cos. We should instead
change to the ajoint registeration mechanism like those in relay.

Followup refactors need to be performed, including:
- Fold the Call.call_type into operator's attribute.
- Enrich the operator registry information
- Refactor passes(e.g. AD, intrin lowering) to use the attribute based transformation

* Fix nms

* Fix remaining testcase

* Address review comment

[COMMUNITY] Matthew Brookhart -> Reviewer (#5886)

keep parameter names from PyTorch (#5887)

Improve type handling in PyTorch frontend (#5834)

* Improve type handling in PyTorch frontend

- Use type information from graph for inputs if available. Check
  against shape information from graph if available.
- Allow user to set default dtype (default to float32 for sanity and
  compatibility).
- Implement type promotion to follow PyTorch mechanism. This includes
  fixing the handling of many "Scalar" overloads in PyTorch binary ops.
- Fix arange/linspace type semantics.
- Added support for traced functions. (Because it really is about the
  "self" input handling.)

Aside from adding an optional default_dtype keyword argument, this does not
change the signature/requirements of from_pytorch.

* Fix scalar detection using numpy.isscalar

and address other review comments. Thank you @siju-samuel

* refine test criteron on qnn_test::test_serialized_modules, fix bool conversion of const

Fail early before running invalid dynamic graphs (#5856)

* fail early before running invalid dynamic graphs

* fix an issue with the VM comment

[Bugfix][Build] Fix building with LLVM-10 on macOS (#5859)

[QUANTIZE] Add nn.batch_flatten as quantizable. (#5805)

* [ONNX] Skip ADD inside Gemm op when vector is zero

* [QUANTIZE] Add nn.batch_flatten as quantizable.

[Target] Introduce Target Id Registry (#5838)

[DOCS] Update has_dtype/has_shape to pattern lang doc (#5847)

Fix map assign issue in CI test (#5854)

Add Python Classes for all Attrs (#5853)

[DataType] Add bfloat16 (#5601)

[Object] Introduce POD-C Compliant tvm::Map (#5740)