Commit Graph

3589 Commits

Author SHA1 Message Date
Siyuan Feng bbc6ba362f [Meta Schedule] Add Auto-Thread Binding Rule (#11177)
The current meta-schedule uses a PostProc `RewriteUnboundBlock` to auto-bind blocks to threads. However, it's a post proc, which means there are no search opportunities, and always splits with `factor=1024`. 

This PR adds a new search rule called `AutoBind` to do a similar thing to bind threads with sampled factors. Also with a corresponding mutator. 

After applying this rule, we get some positive perf results (on RTX-3080):
Element-wise: from 2.76 us to 2.48 us
Conv2d Winograd: from 29.45 us to 18.96 us (ansor 22.00 us)
Resnet18: from  0.591 ms to 0.531 ms (ansor 0.565 ms)
2022-05-20 18:57:37 -07:00
Masahiro Masuda 0274d8e1f1 [TIR] Support tensorization using ldmatrix + MMA (#11355)
* [TIR] Support tensorization using ldmatrix + MMA

commit 3218facf100b0dfc55715acfd1cee156764129ba
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 18 14:04:56 2022 +0900

    some clean up

commit 7a235b69dc2023b3098ed44d591edb63b20a8f4e
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 18 13:55:11 2022 +0900

    parameterize over storage scope in mma store intrin

commit 827ea4c434c35607b241f8e0ae2efe3214ac2458
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 18 13:37:38 2022 +0900

    properly handle floordiv/mod in codegen

commit 42d4c6f42182c9fd79566c0955f99cc82abd5144
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 18 09:53:57 2022 +0900

    update tuned factors for fp16

commit 328d0aa36b2ea9ea1b051970d612bff82d2d20e6
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 18 08:43:30 2022 +0900

    all tests working

commit 5e086cf5fd1404ac38f85c4bfbe692687b45a16c
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 18 07:48:43 2022 +0900

    add doc for mma_fill and mma_store intrin

commit 4f945c4116b6d3bdc965ecb2be2229bb46dc11ab
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 18 06:39:01 2022 +0900

    remove tests

commit df7708f7f67761d9c18f9564bc15abd50c12ac69
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 17 19:52:14 2022 +0900

    unified test

commit 754c83eeb8510b31fb9652b089177f9b8e642ec0
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 17 19:36:24 2022 +0900

    clean up LowerWarpmemory

commit 178c3dcee7bfa17d5d93fec02aa858dc62151670
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 17 19:15:04 2022 +0900

    Use IndexMap

commit 07fb58910338c62847fd902b37801d09b8c673b0
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 17 17:51:44 2022 +0900

    remove 16x8x8 test

commit 2b05b5a5470ac221d559f31a31a8e2ff753b2414
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 17 17:31:35 2022 +0900

    generate mma fill/store

commit bf23fc50f0ffa99e875d9247ca66acec0c36677f
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 17 12:23:30 2022 +0900

    mma intrin generation with meta programming

commit 5afb5f00afd642cb1e39872edc7965f476dcdcb7
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 17 05:26:14 2022 +0900

    ldmatrix intrin generation with meta programming

commit fb62abb3424b88ec48c697e306e05889a3ac306f
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 20:30:49 2022 +0900

    minor

commit 5a80adce24e84d3ec6bf931b60cb9c730d243394
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 19:55:57 2022 +0900

    revert some change

commit e599a55078ee75f2480a721098341812db58cf6f
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 19:54:18 2022 +0900

    remove obsolete files

commit 4b13b85ff91d0d592a7e0c01924e0b49b82f35a8
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 19:51:21 2022 +0900

    wip

commit 848de63455539e25cd0d43e5a65fd048636ef0f7
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 19:44:29 2022 +0900

    wip

commit b35bff97ed10c22559e2164eb7538db0f711ce7e
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 19:31:18 2022 +0900

    update parse error msg

commit ad9b053ef865b1f91f03d7b15ed7aae3420ee213
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 19:26:51 2022 +0900

    fix for avoiding Buffer.vload(...) case

commit 54c686443e370edbfae860d0809b1b6182d26414
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 18:59:55 2022 +0900

    wip

commit 078060fe28d22f1db5f07b1c382dee438f02df60
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 18:57:34 2022 +0900

    wip

commit 576f8415e65e0e8a8a7808885e219b3b53867950
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 18:52:15 2022 +0900

    wip

commit 12a376ae2f44aa6660121e64e0358f2866624f7f
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 17:54:58 2022 +0900

    Squashed commit of the following:

    commit 48eef4981d1a55aaf3b0ac935f2a10347cb1ac2d
    Author: Masahiro Masuda <masahi129@gmail.com>
    Date:   Mon May 16 17:40:48 2022 +0900

        more comment

    commit 8f67fc87038834e9f7e2c5cd3dfe61fabf442206
    Author: Masahiro Masuda <masahi129@gmail.com>
    Date:   Mon May 16 17:11:27 2022 +0900

        update test

    commit ad85036621c005b733763e67ceffae39c356ec99
    Author: Masahiro Masuda <masahi129@gmail.com>
    Date:   Mon May 16 16:54:01 2022 +0900

        add test

    commit 4a5dc3ffd5d0bb4a1700e57897c9e0f26e3d2a88
    Author: Masahiro Masuda <masahi129@gmail.com>
    Date:   Mon May 16 16:40:47 2022 +0900

        [TVMScript] Support function call to help construct AST

commit 76c1bcf0ade45d7433a0066236add8372b1cc547
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Mon May 16 16:30:07 2022 +0900

    simplify iterator in layout transform

commit 936280324ea2c91429a6a85a1b8ee89c7b825928
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Sat May 14 11:31:39 2022 +0900

    remove obsolet files

commit 2e119b422d72d726d5f2bd20fe48a1e62fcb0510
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Sat May 14 10:43:59 2022 +0900

    calculate mma store dst index using inverse affine map

commit 9489434ee52b546e2abb2ab28173eefd51525ba4
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Sat May 14 10:01:12 2022 +0900

    simplify store

commit 1adcb77b8bba8e5d91080fe6cbfc7add7f4365c2
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Sat May 14 09:43:40 2022 +0900

    simplified fill

commit 7b13c736d23e0eac94137aa918101d788e60d4f3
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Sat May 14 09:22:17 2022 +0900

    simplify intrin desc using index map function

commit bcf212dda0f94c51f55c48921f61d92fd3b83777
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Sat May 14 07:16:42 2022 +0900

    seems to work

commit dd8ccf9ec2e48100158152e5d4590d141424e2e2
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Sat May 14 07:11:57 2022 +0900

    poking with the parser

commit 596582cbfbd08ebe23ea71aaf7a447472415ccd1
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 20:04:59 2022 +0900

    16x8x32 4k trans working

commit 273f89a8a6ac34f7c79147563922d34d44bffd08
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 19:52:13 2022 +0900

    add 16x8x16 fp16 trans

commit 8e2066cc4c6e86616bc9751324e63ba81a3b02af
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 19:32:37 2022 +0900

    16x8x16 4k trans working

commit c2d0744051733e94f840d4517bcee9ca5d444c75
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 19:25:52 2022 +0900

    16x8x16 trans working

commit c2e314cdda1c3a931781e51a863901ea178dffec
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 16:19:32 2022 +0900

    tuned int8 4k, 91 TOPS

commit 94d9d965f19ff1a2ebdd342079ef420fb537b16a
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 15:59:33 2022 +0900

    int8 4k tune working

commit 3ca8ca02593aff7540c9655aa831348246171752
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 08:43:57 2022 +0900

    mma 16x8x32 int8 working with ldmatrix b workaround

commit 54f1cb731d4b42a6cbc08baf144e74646400eef5
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 18:23:27 2022 +0900

    wip

commit 9d2844db602dc65af4dbd06a73fdd815f486b8b9
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 16:38:53 2022 +0900

    test tensorize without layout transform

commit 86ee6dabc801aeb8d6917bec6de97b42025dbdd1
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 15:15:34 2022 +0900

    int8 4k tensorize works

commit 39f9e32c9a64222c91daba2c32969b27207a31d2
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Fri May 13 12:44:39 2022 +0900

    begin int8 4k tune

commit 6fa91e55b5ab2ba0f901d0d35be1b2fb3ab092b0
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu May 12 18:53:20 2022 +0900

    try fix ldmatrix b for int8

commit 7a962cddc4799fa3df0c0fdf3c056146d3f2cbdf
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu May 12 18:28:34 2022 +0900

    fixed warp_coeff

commit a0afb5698f307382147a38819e004a2db7f554b1
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu May 12 12:20:01 2022 +0900

    wip

commit f70ccd09b07d5325454ffdc39a7619ea84aa7e06
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu May 12 12:09:57 2022 +0900

    int8 tensorize working

commit 20321fa4674dabc78fe55b5e0e2876c35b245d21
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu May 12 07:06:22 2022 +0900

    starting 16x8x32 int8

commit 441fd193c59cdc436d87ab35896cbb8c779ddf35
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu May 12 05:50:46 2022 +0900

    adding fp16 accum case

commit c9d40b69b1b57bfaddffba09ea07624ae90ee465
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 17:04:29 2022 +0900

    clean up

commit 5b2d48635e762c77c824d1c259ac8bcbcc949421
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 16:38:19 2022 +0900

    16x8x16 4k tune working

commit c3cb170d85600d03da5c3f4cda03552208ca0b8c
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 16:20:27 2022 +0900

    tensoriz fixed

commit 68039b081efcdd6aea1d132940b3745f50164974
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 15:55:25 2022 +0900

    begin 16x8x16 4k tune

commit ced5d8d980cc267d4735957c25cb60d71ae977d2
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 15:50:11 2022 +0900

    16x8x16 worked

commit 3d2c90d77c1bb2df2193e9af6cbaa2bd927a26d8
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 15:47:26 2022 +0900

    fix

commit 403050b03ad6b4f0ee8d45088ffb324727bbae48
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 15:45:10 2022 +0900

    add 16x8x16 test

commit 18e8d73661c99cd1c83021063b41a457afcb1638
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 06:50:32 2022 +0900

    fixed mma store codegen for 16x8x16

commit ec81250561195705122bccb9a2372f71de68121f
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 04:25:25 2022 +0900

    add 16x8x16 mma store codegen

commit e08df2a62a4809bcd39782949283c16e7703aa5c
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 03:47:47 2022 +0900

    tensorized C_warp init

commit ae0678918929c1ceec73f2039467040c5bb7823b
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed May 11 03:06:06 2022 +0900

    mma store codegen working

commit deb4d6646cc93d4cdb4f2560ce723bee4d86e144
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 10 19:22:57 2022 +0900

    update lower warp memory

commit 71fe5fe465300705fa94f9544a2e1a5070de6e0d
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue May 10 09:01:42 2022 +0900

    tensorizing mma store

commit e80a1f148c47f2a3fac2363a733d8d4e2a2631d0
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu Apr 28 19:54:08 2022 +0900

    clean up

commit a9640f4b7c3c9f22b87ca74a61003438dfd8f992
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu Apr 28 19:40:55 2022 +0900

    add tunable 4k test, 36 TFLOPS

commit b9f7eae7041d1a9b3e434c331c874e8347e89dc4
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Thu Apr 28 18:01:08 2022 +0900

    fixed bug in LowerWarpMemory index splitting for ldmatrix

commit 00df30823f874910ed1ec1f74718100311764234
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed Apr 27 07:58:17 2022 +0900

    fixed missing reverse_compute_at

commit 93f9fe7e5f7ad16c8d0e6240c16c0281a0e97dec
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed Apr 27 06:55:12 2022 +0900

    add 4k test

commit 3689ef712aa4b282a4818fa2fa2e7e349c3a5eec
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Wed Apr 27 06:54:09 2022 +0900

    temp disable high dim base indices check in tensorize

commit 0c859c4f385ba0b6f9477b569b80cee80b5b7282
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue Apr 26 19:18:23 2022 +0900

    clean up

commit f6aadbfcfbd73c1667a6de7aedc5894232b8e750
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue Apr 26 19:13:09 2022 +0900

    Add 16x8x8 MMA + LDMatrix test

commit 4cf6b20c6ca415e967ab58d80e4a77c701ad7255
Author: Masahiro Masuda <masahi129@gmail.com>
Date:   Tue Apr 26 18:04:17 2022 +0900

    testing 16x8x8 ldmatrix tensoriation

* set measure_perf to False

* add requires_gpu decorator in tests, always test build on non-ampere

* skip cuda compile on old gpu
2022-05-20 11:33:54 -07:00
Eric Lunderberg 72a5219aad [Schedule] Allowed typing.Tuple in tir.schedule._type_checker (#11289)
* [Schedule] Allowed typing.Tuple in tir.schedule._type_checker

Previously, `typing.Tuple` annotations could not be used with
`tir.schedule._type_checker.type_checked` annotations.  This allows
`Tuple` type annotations to be type-checked.

* Revert change, allow tuples input as List arguments

* Suppress mypy errors

Directly interacting with a type object would otherwise cause some
false positives.

* Corrected unit test for allowing tuples to be used as typing.List

* Represent multi-type lists as List[Union[...]] instead of List[Any]

This gives a better error message and plays nicely with _type2str,
since `typing.Any` doesn't have a `__name__` field.
2022-05-20 10:20:12 -07:00
Andrew Reusch 07d91fa041 Fix function number datatype from char to uint16_t (#11365)
* Fix function number datatype from char to uint16_t

rewrite the modified part to pass lint check

Use 2 bytes for func num in fun_registry

Fix errors in linter

Add the declaration of the helper functions

set 2 bytes for func num in func_registry test units

pass num_func by value

This commit change the datatype of the number of the function from 1 Byte to 2 Bytes.
Besides, I use some helper functions to access the number of function and the first function name.

* Fix aot_executor_module to unbreak CI.

* Fix GraphExecutorModule.

* Remove graph_json_to_c_func_registry.

 * No longer needed and not called anywhere.
 * Superseded by emitting the FuncRegistry directly in codegen.

Co-authored-by: 嚴中璟 <a1245967@gmail.com>
2022-05-20 18:12:55 +09:00
Jiawei Liu a6a34046c4 fix vec*mat in PyTorch converter (#11347)
* fix vec*mat in PyTorch converter

* Trigger CI
2022-05-20 13:45:25 +09:00
Mehrdad Hessar 5e29dddd02 [microTVM][ARM] Add Relay tests for conv2d registered schedules (#11250)
* Added conv2d relay test for each schedule
* Enable relay tests in qemu
* split aot test utils
2022-05-19 16:09:51 -07:00
Farshid Salemi Parizi 8d0da24f12 [Hexagon] moves conftest.py to tvm.contrib.hexagon so outside repos can access the testing fixtures (#11277)
* adding pytest_plugin to python so other repos can access

* import requires_hexagon_toolchain from tvm.contrib.hexagon.pytest_plugin
2022-05-19 09:40:01 -07:00
heliqi ffc0443913 [Frontend] [PaddlePaddle] Add split operator (#11354)
* suuport split op of paddlepaddle

* black formatting
2022-05-19 17:53:03 +09:00
Mehrdad Hessar 9273ea5e49 [Hexagon]Refactor Hexagon_SDK_PATH (#11282)
* refactor HEXAGON_SDK_PATH and remove HEXAGON_GTEST
2022-05-18 13:06:11 -07:00
ah cheng 95509eed26 fix matmul broadcast (#11242) 2022-05-18 10:10:00 -07:00
Elen Kalda 2b1e5ce8dc [microNPU] Fix bug in channels extraction in the matcher (#11335)
* [microNPU] Fix bug in channels extraction in the matcher

If the input tensor layout is in NHCWB16, we were passing W value
instead of the channels to get_valid_block_configs.

* Add test for conv2d
2022-05-18 16:37:05 +01:00
Leandro Nunes f88a10fb00 [TFLite] Add support to int16 data type in TFLite frontend (#10915)
* [TFLite] Add support to int16 data type in TFLite frontend

Add support for int16 data type and int64 biases/accumulators in
the TFLite frontend.

Adjusts TFLite tests to cover int16 convolutions and element-wise;
Fixes a minor typo negtive->negative in the element-wise tests.

* Update src/relay/qnn/op/convolution.cc

Co-authored-by: Elen Kalda <elen.kalda@arm.com>

Co-authored-by: Elen Kalda <elen.kalda@arm.com>
2022-05-18 14:04:24 +01:00
Nicola Lancellotti 1b32245f0b [microNPU] Add a pass to reorder copy and compute nodes (#10959) 2022-05-18 13:22:30 +01:00
Luke Hutton 99caa6533f [TVMC][ETHOSN] Improve target string to avoid duplication (#11272)
* [TVMC][ETHOSN] Improve target string to avoid duplication

Improves the TVMC target string to avoid duplication of the
NPU variant. The new target string will require the just the NPU
name followed by -variant=n78. The old target string is deprecated
and will be removed in a subsequent version of TVM.

Change-Id: I4638f36788df3f478435ac13d3531aad2b23f204

* fix linting

Change-Id: I76a9da511899f24a163be669877605cd1a440022

* fix make variant functions and update test error message

Change-Id: Iff553d4b255c0ce0b86bad42eaa94ee9b1c62508
2022-05-18 12:17:47 +01:00
Masahiro Masuda a4be2ed904 [TVMScript] Support inlined function call as a sugar (#11324)
* [TVMScript] Support function call to help construct AST

* add test

* update test

* more comment

* fix for avoiding Buffer.vload(...) case

* update parse error msg

* wrap func call with try / catch, emit error msg

* silence pylint
2022-05-18 17:56:10 +09:00
Josh Fromm 75c31cae75 [Relay] Bug fix when applying history using an iterator or records. (#11306)
* Bug fix when applying history using an iterator or records.

* I forgot strings are iterables.
2022-05-17 16:17:48 -07:00
czh978 0e2f869eea logsoftmax reusing the softmax function (#11141)
Co-authored-by: caizihua <978497756@qq.com>
2022-05-17 14:08:08 -07:00
Mehrdad Hessar b03f11dfde [Hexagon]Use requires_hexagon instead of requires_hexagon_toolchain if running on hexagon target (#11294)
* refactor requires_hexagon_toolchain

* trigger

* lint
2022-05-17 11:48:04 -05:00
ah cheng 0e27bf5ee9 [frontend][ONNX]support ConvTranspose explicitly specified output_shape (#11076)
* support ConvTranspose explicitly specified output_shape

* fix unit test case

* fix lint test

* retest

* fix code error

* fix lint test

* update test

* retest

* fix test onnx official tests
2022-05-16 09:37:04 -07:00
Ziqang XU 02d57bbc06 [BugFix][Topi] Fix 'duplicated iterator names in the compute definition' bug of roi_align (#11322) 2022-05-16 15:16:30 +09:00
Margaret Qian eed21eeff2 [ONNX] Fix cast op to/from bfloat16 (#11171)
* fix cast from bfloat16

* fix cast to bfloat16 test as well

* clean up comments

* lint

* add comment

Co-authored-by: Margaret Qian <mqian@octoml.ai>
2022-05-15 21:34:37 -07:00
wrongtest 9ab3a1121a [TVMScript] Represent ramp as index slice (#11308)
* support represent ramp as index slice in tvmscript

* fix testcase's comment, check slice lanes instead of extent
2022-05-16 08:33:16 +08:00
Xiyou Zhou 325f6fed56 [MetaSchedule] Allow Easy Logging Level Setting (#11305)
This PR allowed users to set logging level without giving a logger config. Previous implementation hard-coded `logging.INFO` as the default logging level and requires a logger config to change it. Now the logging level and handlers can be inherited from the current `tvm.meta_schedule` logger setting.
2022-05-14 15:19:20 -07:00
Xiaoyu Zhang 87366b56ed Oneflow fronted support more model and fix bug (#11321)
* add relay.f.frontend.fm_oneflow support cnns

* support cuda

* fix mobilenetv2 and reviews

* fix: model without meta info

* support eager and yolo, add test

* fix: license

* add: tutorials

* fix: support new graph

* fix some comments

* refine

* fix concat op convert bug

* refine

* refine

* change cuda to cpu

* fix bug

* fix ci error in tvm

* fix pylint check

* delete useless file

* add skimage package in docker

* fix ci error

* fix bug

* add oneflow fronted test in ci

* merge conflict

* fix tutorial

* try to find error in ci

* revert

* merge conflict

* black oneflow

* Delete from_oneflow.py

* restruct oneflow fronted

* support vision-transformer

* black format

* update black version and reformat

* fix ci error

* fix doc error

* fix gpu fronted test failed

Co-authored-by: hhhfccz <hjk1938927583@163.com>
2022-05-15 04:50:25 +09:00
Alperen Bağ f59c70226b [PYTORCH] [FRONTEND] torch.bool support for data type conversion (#11290)
* [FRONTEND][PYTORCH] Support fo nn.SiLU added

* torch.bool added to torch convert_torch_dtype_map
2022-05-14 18:52:18 +09:00
Mehrdad Hessar aa67a6a01c [Hexagon] Add USMP tests (#11279)
* Add USMP tests

* Address Chris comments

* Address Chris comment on assert

* trigger
2022-05-13 15:38:20 -07:00
Tristan Konolige be65732b2a [ROOFLINE] Roofline analysis over RPC (#11252)
* [ROOFLINE] Roofline analysis over RPC

Run roofline analysis on remote devices if requested. Peak flops and
peak bandwidth estimation are done on the remote device.

* allocate testing arrays directly on device and randomly fill

* forgot to include remote

* lower flops ratio, machine may be using multiple threads

* forgot fill
2022-05-13 14:09:14 -07:00
ibsidorenko d871bbd96f [QNN] Enable constant folding for QNN operations. (#11228)
* [QNN] Enable constant folding for QNN operations.

This commit enables constant folding for QNN operations.
This functionalty is disabled by default, use fold_qnn=True to enable.

Co-authored-by: Alexander Peskov <peskovnn@gmail.com>

* [NFC] Fixed comments

* Added more unit tests for QNN opers in constant folding pass.

* Address PR feedbacks

Co-authored-by: Alexander Peskov <peskovnn@gmail.com>
2022-05-14 04:30:26 +09:00
Eric Lunderberg 0e8107bfc7 [TIR][Arith] Implemented padded inverses in IndexMap (#11235)
* [Debug] Error logging in DetectIterMap

* [Affine] Allowed PrimExpr argument to NormalizeIterMapToExpr

This allows it to be used for any expression containing an
`IterMapExpr`, not just expressions whose top-level node is an
`IterMapExpr`.

* [Affine] Implemented DetectPaddedIterMap

The existing DetectIterMap tries to rewrite index expression as a
linear combination of split/fused iterators, where the new iterators
cover the exact same indices as the original expression.
DetectPaddedIterMap relaxes this condition, allowing the new iterators
to cover a superset of indices that the initial index expression
covered.  It uses the minimum amount of padding necessary to represent
these transformations, and also a predicate that identifies any
padding that has been added.

This is a utility function to be used for layout transformations of
buffers, in cases where the pre-transformation shape of the buffer
does not evenly fit into the post-transformation shape.

* [IndexMap] Implemented IndexMap::NonSurjectiveInverse

Allow non-surjective transformations, with DetectIterMap used to
determine the minimum padding to insert.  Returns the inverse
function, along with a predicate that identifies padding indices.  The
predicate is in terms of the transformed variables.

* [IndexMap] Exposed methods to python

- `IndexMap::Inverse` exposed as `IndexMap.inverse`
- `IndexMap::MapShape` exposed as `IndexMap.map_shape`
- `IndexMap::NonSurjectiveInverse` exposed as `IndexMap.non_surjective_inverse`

* [IndexMap] Extracted _assert_equal_index_map into class method

In preparation for adding additional tests for the IndexMap class,
which will require this functionality.

* [IndexMap] Added unit tests for new behavior

* Re-enabled divisibility check in CheckMapping

Initially disabled as dynamic shapes resulted in padded lengths whose
divisiblity couldn't be proven.  Re-enabled along with a
simplification rule to resolve it.

* Fixed breakage in compute_at primitive

* Corrected typos/examples in docstring
2022-05-13 11:36:50 -07:00
Krzysztof Parzyszek 6c339eaa8f [Hexagon] Remove sim_options from tvm.target.hexagon() (#11293)
We no longer run simulator automatically, so this is not necessary.
Also, the only way to pass options to the simulator was by setting
an environment variable. That variable (HEXAGON_SIM_ARGS) should
be set independently by the user from now on.
2022-05-13 10:07:28 -07:00
Andrey Malyshev c2d1905779 Add Adreno GPU target and topi supporting textures with dynamically allocated textures (#11161)
* Add Adreno GPU target and topi supporting textures

- There are 5 compute/schedules: conv2d for NCHW/NHWC, depthwise_conv2d
  for NCHW/NHWC, average pooling
- Fix of dynamically allocated textures caching
- Add texture-nhwc scope
- Fix issue with codegen of vars having non acceptable symbols

Co-authored-by: Chris Sullivan <csullivan@octoml.ai>
Co-authored-by: Egor Churaev <egor.churaev@gmail.com>

* Address comments

* Add vectorization into some adreno pool flow

Co-authored-by: Li <quic_lih@quicinc.com>

* Fix adreno tests for running on the opencl host platform

* remove unnecessary kDriverVersion in DeviceAttrKind

* Move utils adreno functinos to separate shared file

* fix black hits

Co-authored-by: Chris Sullivan <csullivan@octoml.ai>
Co-authored-by: Egor Churaev <egor.churaev@gmail.com>
Co-authored-by: Li <quic_lih@quicinc.com>
2022-05-13 08:43:06 -07:00
Jacob Bohlin 7c75b77df8 [microNPU] Add various options to the cascader (#10509)
* [microNPU] Added options to Cascader

* Added option to toggle multi-dimensional striping, it is disabled by
  default because it has a very high computational cost. Single
  dimension striping shares most of the benefit with greatly reduced
  cost.
* Added multiple developer/debugging options prefixed with 'dev_'
  Also added these options to tvmc.
* Added cascader logging, if enabled it will dump information about the
  cascader proposals to a 'cascader_log.json' file.

Co-authored-by: Matthew Barrett <matthew.barrett@arm.com>
Change-Id: I2ec59ae0bd84b73b2cc4bc56d39e3831b0aeec27

* Updated memory_reduction testcases

Also added enable_striping to plan_generator.h

Change-Id: I496b30ed6af6f0730087329cd81a69c5040a5e4d

Co-authored-by: Matthew Barrett <matthew.barrett@arm.com>
2022-05-13 15:02:23 +01:00
Mohamad Katanbaf 53fe596682 [rpc] Implemented rpc logging (#11232)
* Implemented rpc logging

* fixing windows build issue

* trigger

Co-authored-by: Mohamad <mkatanbaf@users.noreply.github.com>
2022-05-12 11:48:30 -07:00
Manupa Karunaratne 366a566977 [microNPU] Adding a option to enable striping (#11263)
This commit adds a cascader option to enable
striping explicitly.

When doing so fixed a bug that is associated
with block config selection, that will be
triggered when striping is disabled.

Co-authored-by: Elen Kalda <elen.kalda@arm.com>
2022-05-12 11:31:39 +01:00
Mehrdad Hessar 4eb6497adb [ARM][Strategy] Fix is_int8_hw_support check function (#11193)
* Fix hw schedule condition

* add warning messages to unoptimized schedules
2022-05-11 15:26:48 -07:00
Jiawei Liu 63b6b49f03 fix expand onnx conversion (#11278) 2022-05-12 05:38:21 +09:00
Luke Hutton 3be5622c7e [ETHOSN] Adding support for Leaky ReLU (#11261)
* [ETHOSN] Adding support for Leaky ReLU

Change-Id: Icad69b2ae6ed4b3f3949cf5673efe2571aa66f5f

* add some missing error reporting

Change-Id: I935054c4d19a939e122092fab3c6c77204d9ead8
2022-05-11 12:29:16 +01:00
Luke Hutton cfb5674c25 [ETHOSN] Remove remaining support for the N77 variant (#11262)
Specifically removes some TVMC tests that are no longer necessary
and some partitioning infrastructure.
2022-05-11 09:57:15 +01:00
Matthew Barrett be2ae9433e [TENSORRT] Improvements and fixes for TensorRT (#11203)
A number of small fixes and refactors to improve the robustness of
the TensorRT integration.

Co-authored-by: Mark Shields <mbs@octoml.ai>

Co-authored-by: Mark Shields <mbs@octoml.ai>
2022-05-10 12:39:25 -07:00
Tristan Konolige d2a7f93beb [ROOFLINE] Calculate roofline from existing TIR PrimFunc (#11238)
Refactor roofline_analysis to use a pass instrument to save TIR code
from compilation for feature extraction. This should support different
compilation pipelines and avoids recompiling the module twice.
2022-05-09 14:45:09 -07:00
Mehrdad Hessar 66cc14f69b [Hexagon] Add mobilenet test with AOT (#11204)
* add mobilenet AOT test

* Add _serial_number to super class
2022-05-09 11:28:11 -07:00
Elen Kalda 588679ee83 [microNPU] Add support for conv2d running on two cores on U65 (#10251)
* [microNPU] Add support for conv2d running on two cores on U65

The 512 mac variant has two cores that processes the weights in
parallel, so we need to split the weights and biases into two
and encode them separately.

Change-Id: I53791f614288ac4df181b9462fc632d35b934a86

* Changes due to rebase

* Rebase, improve DivideConstants and expand testing

Make the DivideConstants to operate on non-flattened
tensors to support two core execution in U65.
2022-05-09 16:08:30 +01:00
Christopher Sidebottom 731af42d1b [CMSIS-NN] Increase partitioning accuracy for pooling (#11229)
This ensures that CMSIS-NN is only used when the batch size and layout are correct for the library calls.
2022-05-09 09:23:44 +01:00
Grant Watson e854c0a4d6 [CI] Identify non-inclusive language in commits (#11230)
* Adds a script blocklint.sh that checks for non-inclusive words
  * Updates the task_lint.sh script to call blocklint.sh
  * Replaces the terms Master and Slave where possible
  * Replaces the terms Blacklist and Whitelist
2022-05-09 09:01:38 +01:00
Jian Sheng a3d75ae535 [Frontend][PyTorch] Add: Relay stft operator (#11190)
* Add: Relay stft operator

* fix doc

* address PR comments

* address addtional comments
2022-05-07 15:39:00 +09:00
Wuwei Lin 31be30062b [TIR] Add schedule primitive SetAxisSeparator (#11225)
* [TIR] Add schedule primitive SetAxisSeparator

* remove unused include

* Move ReplaceBufferMutator impl to cc file
2022-05-06 15:21:18 -07:00
Josh Fromm 98aa41e329 [Relay] Flexible shape dispatch transformation (#11199)
* Added pass that creates a semi-dynamic dispatcher around a relay module.

* Added automatic padding feature.

* Output slicing working.

* Multiple input support working i think.

* Added test file.

* Improve comments.

* Fix lint.

* Allow default values.

* Fix docstring.

* Improved documentation based on feedback.

* Add extra check for record loading.

* Improve variable names.

* Add type inference to make sure things worked.

* Added support for multiple outputs.
2022-05-06 12:18:16 -07:00
Leandro Nunes ff7efe767a Revert "Implemented rpc logging (#10967)" (#11227)
This reverts commit aa3bcd9d33, because it
fails on Windows CI as reported in issue #11220. PR #11223 tries to address
it but is is failing in the regular CI with testing issue on Hexagon.
2022-05-06 12:14:20 +01:00
Mohamad Katanbaf aa3bcd9d33 Implemented rpc logging (#10967)
Co-authored-by: Mohamad <mkatanbaf@users.noreply.github.com>
2022-05-04 17:10:18 -07:00
Farshid Salemi Parizi e1acd4b872 [Hexagon] Removes directory after stopping the server (#11212)
* removes hexagon directory
2022-05-04 14:31:14 -07:00