Commit Graph

2692 Commits

Author SHA1 Message Date
Grant Watson f656a2274a Docker env for Arm® Ethos™-U55 Port (#8514)
* Docker env for Arm® Ethos™-U55 Port

* Added Arm® Corstone™-300 Reference System for testing
* Added Arm® Ethos™-U driver stack
* Added installation of Arm® Vela.

Co-authored-by: Manupa Karunaratne <manupa.karunaratne@arm.com>

Change-Id: Ie3cc43943c876d95618a39887aa666da20bcb1e4

* Docker env for Arm® Ethos™-U55 Port

* Removes /opt/arm/cmake/bin from the path
* Parameterizes Arm® Ethos™-U55 driver stack version number

Change-Id: I2162b40f82241fd013643cbfa8847b60d7f4f5a1

* Docker env for Arm® Ethos™-U55 Port

* Adds ethosu as an extra to /python/gen_requirements.py

Change-Id: I2162b40f82241fd013643cbfa8847b60d7f4f5a1

* Docker env for Arm® Ethos™-U55 Port

* Added comment explaining why Vela version needs to be pinned to 2.1.1

Change-Id: I1ade280faa5274cca78899f4dae9e596b16fb5df
2021-08-02 16:06:30 -07:00
Christopher Sidebottom 49d5879e1b Introduce --interface-api={c,packed} parameter (#8280)
* Introduce --interface-api={c,packed} parameter

This introduces structures generated to provide a documented and stable user
friendly interface to a TVM generated model, as can be seen in the AOT
demo application:
```
struct tvmgen_default_inputs inputs = {
  .input_1 = input_data,
};
struct tvmgen_default_outputs outputs = {
  .output = output_data,
};
int ret_val = tvmgen_default_run(&inputs, &outputs, NULL, NULL);
```

To facilitate this, some other changes are included:
* Removed dependency on `aot_executor.{c,h}` in tests, pending the
discussion in the interface RFC as to whether we keep them.
* Moved creation of test DLTensor's into the AOT test utils, in future this
can be replaced by loading via the Python API or otherwise
* Introduce `parametrize_aot_options` which can be used to test
permutations of AOT which work together - for now this filters C
interface and packed operators
* Updated demo application to generate the header for demonstration
purposes, we should consider porting the demo application to Model
Library Format and using the toolchain in the Zephyr App via CMake
instead?

This patch builds upon the improvements @giuseros made to AOT testing
and name mangling from #8014

* Tweak metadata variable description and MLF target loop

* Remove direct usage of `relay::Var` in meta_data.h

This looks like the only place that could be causing the Windows CI failures, so trying removing the additional header in meta_data.h

* Linting fix

* Post-rebase files fixing

These tests were somehow transmuted in transit, I've updated them to the
most recent variant of the test helpers.

* Strip back interface API to just inputs and outputs

This removes any speculative structures from the generated code and cleans up some of the documentation.

* Add header guards and tweak documentation
2021-08-02 16:04:55 -07:00
Junru Shao 7653972954 [Meta Schedule][M3a] Instruction and Trace (#8615) 2021-08-01 18:54:17 -07:00
zhuwenxi 887324f778 [TOPI][CUDA] Improve the performance of scatter_nd (#8479)
* [TOPI][CUDA] Improve the performance of scatter_nd by:

1. Split into 2 kernels, one does the "Init" and another does the "Update".
   Thus they can have different Grid/Block configurations to better utilize
   SMs.
2. Use atomic_add instead of direct assignment, which could avoid the race
   condtion when multiple indices point to the same location of the output
   tensor. With this moidification, it's safe now to use more CUDA threads
   to gain more parallelism.

* Fix python code format.

* FIX: [TOPI][CUDA] Improve the performance of scatter_nd #8479

- Split ScatterND kernel into 2 sub-kernels using ib.new_scope()

- Replace ib.for_range() with blockIdx.y

- Using atomic_add when mode == "add"

- Keep threadIdx.x less than max_threads of GPU

* Comment added

* Add fallback implementation when "mode=add" meets int64

- Atomic_add from CUDA doesn't support int64 data type
- Change "ind{i}" to "ind%d"%i, where names of relay.var could correctly display

* Python format

* Fix line too long

* CI pass

* Empty, for CI pass

* Empty, for CI pass

* Empty, for CI pass

* Empty, for CI pass

* Empty, for CI pass

* Exchange blockIdx.x and blockIdx.y

* check for Vulkan or metal

* Fallback to previous algorithm when mode==update

* Update python/tvm/topi/cuda/scatter.py

Co-authored-by: Tristan Konolige <tristan.konolige@gmail.com>

* Assign TODO

* Swapping then and else block

Co-authored-by: wenxizhu <wenxizhu@tencent.com>
Co-authored-by: CaptainDuke <captainduke328@gmail.com>
Co-authored-by: Tristan Konolige <tristan.konolige@gmail.com>
2021-08-01 10:35:43 +09:00
masahi 4b67daccb9 [CUDA] Support multiple TIR-level dynamic shared memory allocations (#8571) 2021-07-31 10:20:24 -04:00
Tianqi Chen 7d8a774a6b [VTA] Recover rpc server support (#8604) 2021-07-31 08:38:08 -04:00
Siyuan Feng 2a8950b5c0 [TensorIR] Support for match_buffer from subregion (#8585)
Co-authored-by: Junru Shao <junrushao1994@gmail.com>
Co-authored-by: Bohan Hou <32121147+spectrometerHBH@users.noreply.github.com>
Co-authored-by: Ruihang Lai <lairuihangdongdong@qq.com>
Co-authored-by: Hongyi Jin <3231950289@qq.com>
Co-authored-by: Wuwei Lin <wuwei@apache.org>
2021-07-30 22:06:38 -07:00
Ruihang Lai 5012462ef8 [TensorIR][M2a] Reduction Factoring (RFactor) (#8544)
Co-authored-by: Junru Shao <junrushao1994@gmail.com>
Co-authored-by: Siyuan Feng <Hzfengsy@sjtu.edu.cn>
Co-authored-by: Bohan Hou <32121147+spectrometerHBH@users.noreply.github.com>
Co-authored-by: Hongyi Jin <3231950289@qq.com>
Co-authored-by: Wuwei Lin <wuwei@apache.org>
2021-07-30 22:04:02 -07:00
Matthew Brookhart 22c7d6107f speed up reference resize kernel (#8592) 2021-07-30 09:49:19 -04:00
Altan Haan 81480287a8 [FIX][CI] hotfix check_grad perf regression (#8581)
* hotfix check_grad perf regression: lift compile out of hot loop

* hoist interpreter creation out of python closure, fix weird conv2d bug on arm cpu

* lint

* try one more fix
2021-07-30 09:09:37 +01:00
Chenfan 88dd31b600 [VM] Bug fix for numpy scalar input in vm (#8553)
* Bug fix for numpy scalar input in vm

* Bug fix

* Re-triggle CI

* Update

* Update UT

* Re-triggle CI
2021-07-30 14:11:51 +09:00
masahi bef7bf9b2b [Refactor] Remove AttrStmt with storage_scope key (#8516)
* Remove all attr::storage_scope usage

* pyformat

* fixed VTA tests

* Update TIR text printer to print storage_scope on allocate

* print storage scope in AllocateNode ReprPrinter

* Fixed accidently removed scope tag check

* remove unused function

Co-authored-by: masa <masa@pop-os.localdomain>
2021-07-29 14:29:14 -05:00
Chenfan 850abb0c01 [TOPI] Add transpose_a/b & dynamic shape support for batch matmul (#8527)
* Add basic support for batch matmul transpose

* Update

* Lint fix & add tf convert support

* Update

Lint fix

* Bug fix for qnn.batch_matmul

* Bug fix for tensorflow test

* Add grad support for batch_matmul

* Lint fix

Re-triggle CI

Bug fix

Re-triggle CI

Re-triggle CI

Re-triggle CI
2021-07-29 10:15:21 -07:00
Mario Perić SiMa.ai 83ce7fe827 [TOPI] Fix nn.pool*d issue with 'vectorize' function and add unit tests (#8541)
* Fix issue in 'vectorize' function for 1D and 3D tensors

* Add pooling tests for channel last layouts

* Add support for more general layouts in "poolnd" implementation

* Reformat with 'black'

* Fix lint issues
2021-07-29 19:53:02 +09:00
zotanika a9933211a9 [TOPI][CUDA] minor change on assert statement in conv2d_NCHWc_int8.cuda (#8554)
* [TOPI][CUDA] minor change on assert statement

* [TOPI][CUDA] reformatting
2021-07-29 19:33:44 +09:00
Altan Haan b6b8269d71 add support for softmax and log_softmax with MIOpen (#8543) 2021-07-29 18:36:20 +09:00
Cody Yu 74715f4459 [AutoScheduler] Fix task extraction with TE compiler (#8560)
* [AutoScheduler] Fix task extraction with TE compiler

* fix

* test

* Update python/tvm/auto_scheduler/relay_integration.py
2021-07-29 15:18:38 +08:00
AndrewZhaoLuo 92f9558e90 [Coreml] Fix Coreml Input Shape Handling (#8562)
* convert ot python list like expected

* test example

* jostle ci

Co-authored-by: Andrew Zhao Luo <andrewzhaoluo@system76-pc.localdomain>
2021-07-29 15:13:40 +08:00
Lunderberg 6b2cbfe31f [Contrib] Added default non-verbose to download_testdata(), pass to download() (#8533)
* [Contrib] Added default non-verbose to download_testdata(), pass to download().

Minor cleanup as well, while in the file

- Using tempfile.TemporaryDirectory instead of explicit cleanup.

- Pass through verbose/retries arguments if replacing a corrupted
  copy.

* [Contrib] Switched download.py from print statements to logging

* [Contrib] Added shutil.copy2 fallback after downloading file.

Initial implementation using tempfile.TemporaryDirectory assumed that
the tempdir and output location were on the same drive, and could be
renamed.  This update falls back to copying from the temporary
directory, in case the tempdir is on a different drive.

Co-authored-by: Eric Lunderberg <elunderberg@octoml.ai>
2021-07-29 15:35:12 +09:00
pzq a17ee9f54a [BUGFIX] fix illegal memory access bug in reduce op schedule by constriant threadIdx.y (#8566)
Signed-off-by: ziqiang.pzq <ziqiang.pzq@alibaba-inc.com>

Co-authored-by: ziqiang.pzq <ziqiang.pzq@alibaba-inc.com>
2021-07-28 15:39:33 -04:00
Rohan Mukherjee 8d9df2dd5d [TensorRT, BYOC] Handling a corner case in TRT RemoveDropout pass (#8506)
* [TensorRT, BYOC] Handling a corner case in TRT RemoveDropout pass

* changing visit logic
2021-07-28 08:36:35 -04:00
Xingyu Zhou 496f10b007 [Frontend][TENSORFLOW] Add support for unpack with dim 0 after tensorlist stack (#8558)
* enable testcase when tensorlist stack follows by a unpack for dim 0

* address reviews and improve the docstring
2021-07-27 21:20:58 -04:00
Tianqi Chen 3b7aed3263 [FFI] Specifically check handle for recursion during shutdown (#8548)
NOTE: previously slot may get overriden by child class and it
is better to directly check for handle here.
2021-07-27 08:52:49 -04:00
Joey Chou SiMa.ai a492db89d9 [Bugfix] Visit each input param of the function in ExprVisitor visit_function (#8521) 2021-07-26 17:23:45 -07:00
Josh Fromm ee207fd796 [RPC] Add explicit type cast to print. (#8524) 2021-07-26 19:48:40 -04:00
Elen Kalda 9c63f4fc31 [TFLite] Mimic the TFLite's 2.4 reader's behaviour (#8538)
In TFLite 2.4, the builtin code value can be either in
"deprecated_builtin_code" field or "builtin_code" field (as long
as the value is less than 127) and similarly to the TFLite's
reader, we should use the higher value of the two.

Change-Id: I0d738f9257187903b4c5b4cc5a8733a451ddc02e
2021-07-26 12:17:57 +01:00
Xingyu Zhou 8ab2074944 [Frontend, Tensorflow2] Added support for TensorList ops (#8454) 2021-07-25 00:54:22 -07:00
Tianqi Chen e664ef0be8 [PRINTER] Fix the repeatitive cast in scripr printing (#8531) 2021-07-24 16:53:00 +09:00
Tristan Konolige bce0db55d7 [PROFILING] Add json output to profiling reports (#8503)
* [PROFILING] Add json output to profiling reports

* format json in comments
2021-07-24 15:57:40 +09:00
Mehrdad Hessar 59e96e04ce [microTVM][Cortex-R5] Add zephyr cortex-r5 board to Zephyr (#8519)
* cortex r5 added

* add aot demo
2021-07-22 09:33:19 +01:00
Lunderberg 45497bd362 [Topi][UnitTests] Parameterize conv2d and depthwise_conv2d tests (#8433)
* [UnitTests][Topi] Updated test_topi_conv2d_nchw.py to have parametrized tests.

- Better error messages, displays which workloads/targets failed and why.

- Fixed bug in topi.nn.conv2d._get_workload exposed by the
  parametrized tests.  Incorrect padding if the "SAME" parameter is
  used with dilation>1.

- Fixed bug in tvm.topi.x86.group_conv2d._get_default_config, missing
  dilation parameter in call to _get_conv2d_workload.

* [UnitTests][Topi] Parametrized the tests in test_topi_depthwise_conv2d.py

In preparation for parametrizing to test on float16 as well.

- Single test_conv2d test with parameters for layout/input sizes.

- Extended the support for NCHWc layouts, so that they could be
  included in the parametrization.  (Implemented
  topi.testing.depthwise_conv2d_python_nchwc and
  topi.nn.scale_shift_nchwc, added layout argument to
  topi.nn.depthwise_conv2d._get_workload).

Co-authored-by: Eric Lunderberg <elunderberg@octoml.ai>
2021-07-22 10:26:08 +09:00
schilkunda-amba 789ab1f406 [Relay] Support resize in the ONNX conversion (#8455)
* [Relay to Onnx]

* Added support for resize2d op
* Added unit test

* [Relay to Onnx][Resize]

* Fixed formatting errors

* [Relay to Onnx][Resize]

* Fixed issue in resize conversion: round maps to round_preferc_ceil
* Updated resize unit test to test for coordinate transform mode and
round
* Known issue: Does not match for (NN, align_corners) and Cubic

* * Fixed formatting errors

* * Fixed some more formatting errors
2021-07-21 13:59:21 -07:00
Hongyi Jin 2df0854a2b [TensorIR][M2a] Fuse, Split (#8467)
* Fuse&split (#408)



Co-authored-by: jinhongyi <323195289@qq.com>
Co-authored-by: Junru Shao <junrushao1994@gmail.com>
2021-07-21 16:59:06 -04:00
ziheng eacc2cb0d3 [TIR] Bugfix for zero number arguments tir functions. (#8515)
* [TIR] Bugfix for zero number arguments tir functions.


Co-authored-by: Junru Shao <junrushao1994@gmail.com>
2021-07-21 12:20:59 -04:00
masahi 1a1be09c24 [Refactor] Remove scope attribute from Buffer class (#8463)
Co-authored-by: masa <masa@pop-os.localdomain>
2021-07-20 16:19:46 -04:00
Andrew Reusch ca28dfff35 Rename runtime-config to executor-config and add documentation for Model Library Format (#8270)
* Rename runtime-config to executor-config.

* Add documentation.

* address comments, make tests pass

* fix unit test

* fix sphinx doc errors

* address manupa comments
2021-07-20 20:14:09 +01:00
Valery Chernov d5818b4cd3 [Frontend, pytorch] Vc/pytorch lstm (#8447)
* lstm layer conversion to relay from pytorch model (TorchScript) was supported

* bidirectional LSTM layer was supported for pytorch API

* lstm tests were implemented. fixes in pytorch lstm

* fix pytorch bidirectional lstm. update test comment

* black format and some small fixes

* LSTM with projection was supported for pytorch frontend. test was updated by new combination of LSTM types

* lint fixes

* add bias switcher for LSTM types test. fix LSTM implementation in pytorch frontend for case without biases. exception in the test for conversion LSTM with projection from pytorch to ONNX

* transfer test_lstms to pytest format

* onnx model saving was implemented through io.BytesIO. creating/removing tmp dir was removed. remove unneccessary comments

* gpu target was added to the test

Co-authored-by: Valery Chernov <valery.chernov@deelvin.com>
2021-07-20 21:36:57 +09:00
Matthew Brookhart 6d88bdd3ea Update Docker CI (#8193)
* add failing onnx tets

* point jenkins at new docker

* support convtranspose opset 11 autopadding

* Don't force output shape for conv transpose tests, add 1D and 3D cases

* disable test until CI update complete

* try updating docker images again

* skip a test until update complete

* next try at docker images

* manage TF memory use in TF1 tests

* support explicit padding for NCHW TF padding test

* Update to tagged tlcpack images

Thanks, Andrew!

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

Co-authored-by: Andrew Reusch <areusch@gmail.com>
2021-07-19 09:18:49 -06:00
Rijul Gupta f64ddff497 Add missing shape functions for relay.nn operations (#8489)
* Update _nn.py

add a few missing shape functions

* Update _nn.py

Updated conv_transpose shape function to accomodate conv1d_transpose

* added tests for new functions

* fixed a lint error

* fixed shape func error

* attempt fixing cuda error
2021-07-19 10:39:29 +08:00
mbs-octoml 2c61afa4d1 Switch from CompileEngine to TECompiler in Interpreter (#8486)
This continues on:
https://discuss.tvm.apache.org/t/rfc-relay-tecompiler-rewrite-existing-compile-engine-to-match-updated-compiler-flow/9233
and #751, this time just replacing CompileEngine with TECompiler in the Interpreter,
using the JIT helper added to help the transition.

Some whitespace improvements while there.
2021-07-16 23:57:18 -07:00
Huang, Guangtai 0545962002 [Bugfix] [tir] do not simplify 'Any() - Any()' to 0 (#8266)
* fix

* fix lint

* remove

* address comments
2021-07-16 23:56:18 -07:00
Tantalus13A98B5F 7456cfc531 [AutoTVM] Re-enable ref_input (#8113)
* [AutoTVM] Re-enable ref_input

* add ref_input on measure_option

* add ref_input unittest

* fix: test reformat

* [autotvm] [ref-input] refine test and description

* [autotvm] [ref-input] revert arg on measure_option
2021-07-16 23:54:56 -07:00
srinidhigoud c95d16e097 [Frontend][Tensorflow2] Stridedslice and concat_v2 fix (#8483)
* fix for strided_slice when begin > end in case of shrinkaxis_mask

* fix for name_hint missing error for concat_v2 op

* removing a local fix

* adding more testing capability to concat_v2
2021-07-17 15:26:41 +09:00
Honghua Cao 44f1baf794 [TVMSCRIPT] Support tir.abs node in tvm script (#8488)
Co-authored-by: honghua.cao <honghua.cao@streamcomputing.com>
2021-07-16 22:13:58 -04:00
Alexander Pivovarov 7388715632 Fix _get_yolo_detections (#8477) 2021-07-16 10:26:25 -07:00
Chenfan cba9cf3d22 [VM] Fix the shape function of conv nhwc (#8480)
* Add dynamic support for conv2d nhwc
2021-07-16 20:22:36 +08:00
Masahiro Hiramori 2b57907d9e fix minor misspelling (#8476)
Co-authored-by: Masahiro Hiramori <mshr-h@users.noreply.github.com>
2021-07-16 09:57:27 +01:00
Y c263f22ffd [TVMC][FIX] Compiler supports input with a slash (#8481) 2021-07-15 23:08:58 -07:00
Raghav-Chakravarthy 5bb01ef8ea fix typo (#8484)
* fix typo

* Fixed typos in documentation
2021-07-15 23:08:10 -07:00
Josh Fromm ce15ca65ef [Relay][Frontend][ONNX] Allow importing models with malformed Loop nodes. (#8475)
* Snapshot

* Undo comments.

* Add testing for malformed loop nodes.

* Format oops.
2021-07-15 12:13:21 -06:00