Commit Graph

3931 Commits

Author SHA1 Message Date
Luke Hutton 7c318d7196 [ETHOSN] Remove support for older versions of the driver stack (#12347)
Removes support for driver stack versions older than 22.05
(semantic 3.0.1). Additionally, changes the integration to make
version checks using semantic versioning rather than the previous
year.month versioning method.
2022-08-22 11:31:44 +01:00
chengven027-intellif e9aad35cf3 fix group conv3d pack kernel shape error (#12523) 2022-08-22 17:01:23 +09:00
Lite Ye 262906516a [TVMScript] Printer: add boolean operators to OperationDoc (#12518)
This PR adds boolean operators to OperationDoc. This is needed by the TIR expression printing because it has `tir::And` and `tir::Or`.

Tracking issue: #11912
2022-08-21 22:10:23 -07:00
Lite Ye cc769fdc95 [TVMScript] Printer entry point (#12462)
This PR:

- Adds an entry point for the TVMScript Unified Printer
- Adds a helper object class `RootNodeContainer` to provide an injection point for the actual printer implementation to add specialized logic on the root node to print.

Tracking issue: https://github.com/apache/tvm/issues/11912
2022-08-20 20:41:42 -07:00
Junru Shao 8ee4b604bc [MetaSchedule] Migrate MemoryDatabase to C++ (#12514)
This PR migrates the existing MemoryDatabase, which is implemented in
python at the moment, to C++. The original intent of having an in-memory
database that does not persist on disk is merely for testing, but as
times go on, we found it useful in production workflow, and thus decided
to migrate it C++ for potentially better performance.
2022-08-20 18:11:58 -04:00
Junru Shao 125c9caa02 [MetaSchedule] Implement ScheduleFn as a C++ class (#12513) 2022-08-20 18:11:13 -04:00
Junru Shao 3b3443bd12 [TIR][Schedule][UX] Beautify TIR Trace Printing (#12507)
Following https://github.com/apache/tvm/pull/12197, this PR introduces
`Schedule.show()` which convenience the user experience in the following
two aspects:
- Python syntax highlighting
- Outputs a schedule function instead of standalone instructions so that
it's easier to follow.

To demonstrate this change:
- Before `Schedule.show()` is introduced:
<img width="555" alt="image" src="https://user-images.githubusercontent.com/22515877/185713487-03722566-1df7-45c7-a034-c1460d399681.png">

- After this change:
<img width="583" alt="image" src="https://user-images.githubusercontent.com/22515877/185713564-c54f3a9d-cd52-4709-a8b8-d8a61361e611.png">
2022-08-20 18:10:19 -04:00
Mehrdad Hessar c83ee08c10 fix pytest (#12483) 2022-08-19 16:35:14 -07:00
Yulv-git bdcfa01eae [Fix] Fix some typos (#11503)
Fix some typos in src/.

Co-authored-by: driazati <driazati@users.noreply.github.com>
2022-08-19 16:31:30 -07:00
chengven027-intellif 9d6039b879 fix group_conv3d caculate error (#12500) 2022-08-20 08:30:47 +09:00
Mehrdad Hessar 8b3401ce6b [microTVM] Add config space to dense_dsp schedule (#12444)
* add config space

* lint

* lint
2022-08-19 15:40:08 -07:00
Krzysztof Parzyszek 6def53aeaa [Target] Only append default keys if target doesn't have any yet (#12474)
* [Target] Only append default keys if target doesn't have any yet

This allows target parsers to provide their own target keys. Without this
change, the default keys would always be appended, which may or may not
be desirable.

* Add "cpu" to ARM CPU keys

* Add "cpu" to the keys in the mprofile target parser

* Restore the mprofile cpptest, since the "cpu" key is back

* So the -device attribute is actually needed...
2022-08-18 18:11:23 -05:00
Aakanksha Verma 88928a40f1 [HEXAGON] Auto-vectorization (fp16) for v68 (#12397)
* Auto-vectorization (fp16) for v68

* use tvm.testing.main in fp16 test of tanh_slice op
2022-08-18 12:40:05 -05:00
Yaxing Cai 250b68e202 [TVMScript] IRBuilder, IRBuilderFrame base class (#12482)
* [TVMScript] IRBuilder, IRBuilderFrame base class

This PR introduces basic data structures of the generic IRBuilder
across the codebase.

IRBuilder is a general-purpose IRBuilder that can be used in TIR, Relax
and any other vendor-specific dialects; IRBuilderFrame is where contexual
information as stored in the IRBuilder.

* fix linter

* Update include/tvm/script/ir_builder/base.h

Co-authored-by: Junru Shao <junrushao1994@gmail.com>
2022-08-18 20:43:00 +08:00
arangasa 436c17f885 [HEXAGON][TOPI] This PR adjusts schedules so >64 length vector loads/stores are not generated at LLVM level. This is a workaround for an instruction selection issue in current version of llvm for hexagon (#12471) 2022-08-18 17:55:44 +09:00
Black f64a3bda25 [Frontend][Pytorch] Add axis N when maxpool3d layout is (C,D,H,W) (#12467)
* Add axis N if input is (C,D,H,W) layout.

* Add (C,D,H,W) test case.
2022-08-18 09:22:42 +09:00
Yaxing Cai b0b9bd976b [TIR] Expose TVM Backend API-related Builtins and Misc (#12468)
This PR exposes the following TIR operation in python:

`tvm_thread_allreduce`: tested [here](https://github.com/apache/tvm/blob/bcc7cde95c1e84b85f18c07110489350865b8cfe/tests/python/unittest/test_tvmscript_type.py#L135)
`type_annotation`: tested [here](https://github.com/apache/tvm/blob/bcc7cde95c1e84b85f18c07110489350865b8cfe/tests/python/unittest/test_tvmscript_roundtrip.py#L718)
`tvm_access_ptr`: tested [here](https://github.com/apache/tvm/blob/bcc7cde95c1e84b85f18c07110489350865b8cfe/tests/python/unittest/test_tvmscript_roundtrip.py#L717)
`tvm_throw_last_error`: tested [here](https://github.com/apache/tvm/blob/bcc7cde95c1e84b85f18c07110489350865b8cfe/tests/python/unittest/test_tvmscript_roundtrip.py#L343)
`TVMBackendAllocWorkspace`: tested [here](https://github.com/apache/tvm/blob/bcc7cde95c1e84b85f18c07110489350865b8cfe/tests/python/unittest/test_tvmscript_roundtrip.py#L340)
`TVMBackendAllocWorkspace`: tested [here](https://github.com/apache/tvm/blob/bcc7cde95c1e84b85f18c07110489350865b8cfe/tests/python/unittest/test_tvmscript_roundtrip.py#L465)

Co-Authored-By: yongwww <yongcale@gmail.com>
2022-08-17 12:06:19 -07:00
Yaxing Cai bcc7cde95c Expose Missing TIR Builtins to Python (#12466)
This PR exposes the following TIR operation in python:

`address_of`: tested [here](https://github.com/apache/tvm/blob/d2f9f254d275df256dbcbc5a9f8b3a07cee1d81f/tests/python/unittest/test_tvmscript_roundtrip.py#L3247)
`lookup_param`: tested [here](https://github.com/apache/tvm/blob/d2f9f254d275df256dbcbc5a9f8b3a07cee1d81f/tests/python/unittest/test_tir_usmp_analysis_extract_bufferinfo.py#L171)
`infinity`: add new unittest
`reinterpret`: tested [here](https://github.com/apache/tvm/blob/d2f9f254d275df256dbcbc5a9f8b3a07cee1d81f/tests/python/unittest/test_tvmscript_roundtrip.py#L2991)
`isnullptr`: tested [here](https://github.com/apache/tvm/blob/d2f9f254d275df256dbcbc5a9f8b3a07cee1d81f/tests/python/unittest/test_tvmscript_roundtrip.py#L260)

Co-Authored-By: yongwww <yongcale@gmail.com>
2022-08-17 02:26:45 -07:00
Yaoda Zhou 073304dadb [TVM PyTorch Integration] libstdc++ CXX11 ABI Compatibility & boolean tensor support (#12232)
* first commit

* rename

* cmake

* deprecated

* newline

* config

* config

* typo

* skip tvm_class

* rename

* delete ptr

* delete ptr

* save progress

* boolean support

* cmake file

* polish code

* compile config

* improving the codes

* format

* doc&errormsg

* zero-cost copy

* one step

* to ndarray

* extra output

* delete extra codes

* update test

* boolean support

* strong test

* decrease memory copy

* polish

* reformat

* polish

* remove redundant import

Co-authored-by: juda <yzhou@octoml.ai>
2022-08-17 17:33:37 +09:00
Krzysztof Parzyszek d2f9f254d2 Use std::string_view, remove experimental or pre-14 variants, NFC (#12460) 2022-08-16 21:38:29 -07:00
Lite Ye 1a9faeb77e [TVMScript] Printer IRDocsifier (#12396)
This PR:

- Adds IRDocsifier

This PR is in draft state because it's branched off from a pending PR #12336

Tracking issue: https://github.com/apache/tvm/issues/11912

Co-authored-by: Greg Bonik <gbonik@octoml.ai>
2022-08-16 14:06:29 -07:00
Yaxing Cai b831b06f36 Expose Struct/Tuple-related TVM Builtins (#12452)
This PR exposes the following TIR operation in python:

`tvm_tuple`: tested [here](https://github.com/apache/tvm/blob/c477c763c37adf29b34528ca52d231d622719b3e/tests/python/unittest/test_tvmscript_roundtrip.py#L554)
`tvm_struct_get`: tested [here](https://github.com/apache/tvm/blob/c477c763c37adf29b34528ca52d231d622719b3e/tests/python/unittest/test_tvmscript_roundtrip.py#L200)
`tvm_struct_set`: tested [here](https://github.com/apache/tvm/blob/c477c763c37adf29b34528ca52d231d622719b3e/tests/python/unittest/test_tvmscript_roundtrip.py#L2432)

Co-Authored-By: yongwww <[yongcale@gmail.com](mailto:yongcale@gmail.com)>

cc @junrushao1994
2022-08-16 11:58:04 -07:00
Nicola Lancellotti 09a4ac48ed [ETHOSN] Add support for Requantize (#12384)
This commit adds support for the requantize operator for the Arm(R) Ethos(TM)-N NPU.
2022-08-16 10:17:17 +01:00
Yaxing Cai c477c763c3 [TIR] Expose Misc TIR operations to python (#12435)
This PR exposes the following TIR operation in python:
- `assume`: tested [here](https://github.com/apache/tvm/blob/bb513866ad70fa20eb0c37ca339d330d6a76c747/tests/python/unittest/test_tir_transform_remove_assume.py#L34)
- `undef`: tested [here](https://github.com/apache/tvm/blob/bb513866ad70fa20eb0c37ca339d330d6a76c747/tests/python/unittest/test_tir_transform_remove_undef.py#L63)
- `likely`: tested [here](https://github.com/apache/tvm/blob/bb513866ad70fa20eb0c37ca339d330d6a76c747/tests/python/unittest/test_tir_schedule_compute_at.py#L849)

Co-Authored-By: yongwww <yongcale@gmail.com>
2022-08-16 00:35:31 -07:00
Siyuan Feng 55f1d7e2c4 [TIR][UX] allow override when register TensorIntrin (#12439)
* allow override when register TensorIntrin

* lint
2022-08-15 10:02:35 -07:00
Yaxing Cai bb513866ad [TIR] Expose Stack-related TVM builtins in Python (#12429)
Added the following operations in TIR:
- `tvm_stack_alloca`
- `tvm_stack_make_shape`
- `tvm_stack_make_array`

Co-Authored-By: yongwww <yongcale@gmail.com>
2022-08-14 22:01:13 -07:00
Yaxing Cai d805ae3bd9 [TIR] Expose: call_packed_lowered, call_cpacked_lowered (#12425)
Added the following operations in TIR:
- call_packed_lowered
- call_cpacked_lowered

Co-Authored-By: yongwww <yongcale@gmail.com>
2022-08-14 21:58:52 -07:00
Yaxing Cai e111984117 [TIR] Allow tir.Buffer converted to BufferLoad/BufferRegion with __getitem__ (#12422) 2022-08-14 07:43:58 -07:00
Yaxing Cai e3d0305881 [TIR] Avoid import * in TIR tensor intrinsic registration (#12424) 2022-08-14 07:43:38 -07:00
Yaxing Cai 57a3b26343 [Fix] Fix dtype in Cache-Read/Write (#12421) 2022-08-14 07:43:21 -07:00
AndrewZhaoLuo 8ec7a0a460 [TIR] Fix assert for tensorcore int8 intrinsics (#12365) 2022-08-13 23:15:34 -07:00
Yaxing Cai 0ac5dd7f64 [Fix] Fix errors in error checking and reporting (#12423) 2022-08-13 20:07:45 -07:00
Huan Mei d8846ec975 [TOPI][OP]change float multiplication of resize op to integer division (#12315)
* [TOPI][OP]change float multiplication of resize op to integer division

* add unittest.

* 1. add docstring
2. rename param
2022-08-14 10:39:56 +08:00
Lite Ye d33a332283 [TVMScript] Printer VarTable (#12336)
This PR:

- Adds VarTable for the new TVMScript Printer

Compared to the prototype version, this:

- Removes unnecessary public methods.
  - GetObjectName
  - GetUniqueName
- Add Frame parameter for `Define` methods. VarTable will add callback to Frame to remove variable when Frame exits.
- Changes DocFactory from `ExprDoc(ObjectPath)` to `ExprDoc()` to simplify var definition.

Tracking issue: https://github.com/apache/tvm/issues/11912
2022-08-13 13:04:39 -07:00
Wuwei Lin 036aa722ae [TIR] Add pass ManifestSharedMemoryLocalStage (#12355)
Added a pass to insert local (cache) stage for the shared memory. It's similar to cache read but bypasses the limitation of int set analysis for compacting buffer region by inferring the buffer shape from the loop extents.
2022-08-13 01:06:50 -07:00
Wuwei Lin 96cac7501d [MetaSchedule] Filter vector_load_lens based on buffer dtype (#12408)
This makes the same config generic to work across workloads with different types.
2022-08-13 01:06:16 -07:00
Alexey Voronov 779a7ada4f Add sort_by_time flag to debug_executor.run method (#12402) 2022-08-12 11:15:45 -07:00
Egor Churaev c3c7c4ccc3 [Profiler] Fix graph_executor_debug hang (#12382)
For some operations such as `__nop` or `__copy` the measured inference
time is equal to 0. In this case we are in infinite loop and we won't
exit from it. Added new parameter `limit_zero_time_iterations ` which specify the
maximum number of repeats then the inference time is equal to 0. When
we exceed this value then we will exit from a loop.
2022-08-12 09:26:38 -07: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
Elen Kalda d874a8ed52 [microNPU] Reorder copies and computes based on the cycle count (#11591)
If the cascader is enabled and the ops in TIR have the cycle
count annotation, enabling the reorder_by_cycles option will
reorder to copies and computes based on a cycle count.

If reorder_by_cycles is enabled, max_copy_movements is ignored.

This pass is currently not part of the TIR pipeline since it
assumes that weights and bias of a compute op are merged into
one constant (which is WIP).
2022-08-12 16:20:08 +01:00
Yuanjing Shi a1c371f46c [PyTorch] Fix all_any_common with no default input (#12395)
* fix all_any_common with no default input

* work around

* better naming
2022-08-12 18:09:22 +09:00
Yuanjing Shi 22dcf4490d [PyTorch] Fix pad_common for float pad_value (#12134)
* fix pad

* fix constant padding and handle float infinity

* revert change to pad_width

* fix constant pad value
2022-08-12 16:02:48 +09:00
Lite Ye e3e65ce581 Add Python function to get type index by class (#12393) 2022-08-11 20:36:47 -07:00
Lite Ye 3cb57d20a0 [TVMScript] Printer Frame (#12366)
This PR:

- Implement Frame for the TVMScript Unified Printer

Compared to the prototype version, this:

- Removes the dependency of VarTable (SymbolTable) from Frame
- Adds a callback array to the Frame base class so that VarTable can add callback to clean variable when Frame goes out scope

Tracking issue: https://github.com/apache/tvm/issues/11912
2022-08-11 20:22:37 -07:00
Egor Churaev 5deb95a947 [Adreno][OpenCL] Get rid of extra memory copy (#12286)
* Add annotation pass for device_copy where we get buffers but expect

textures

* Fix issues with running device_copy

* Get rid of extra memory copy

* Fix build after cherry-picking

* Fix lint

* Fix CI

* Apply comments

Co-authored-by: Andrey Malyshev <elvin.nnov@gmail.com>
2022-08-12 11:53:39 +09:00
Krzysztof Parzyszek 99f5e921a5 Update C++ standard to C++17 (#12337)
* Update C++ standard to C++17

LLVM has switched to C++17 in its development branch. Follow suit to be
able to compile LLVM headers.

* Clang 8.0+ also supports -faligned-new

* Make make verbose

* Use CMAKE_OSX_DEPLOYMENT_TARGET to set minimum macOS version (10.12)

* Update llvmdev version in conda to >= 11

Something seems to be wrong with the llvmdev 10.0.0 packages, since the
LLVM unit test fails on Windows. It works fine when LLVM 10 is compiled
from sources.
2022-08-11 17:21:53 -05:00
Farshid Salemi Parizi 48354ded38 [Hexagon] Add skip option for RPC server initialization (#12368)
* add hardware parallelism in hexagon

* fix name

* lint issue

* better description
2022-08-11 12:23:06 -07:00
Florin Blanaru ecfd9692a0 Unify name mangling in TVM (#12066)
* Add NameSupply and GlobalVarSupply

* Build GlobalVarSupply from IRModules instead of having it attached to an IRModule.

* Pass GlobalVarSupply when lowering shape funcs

* Partially replace instantiations of GlobalVar with GlobalVarSupply

* Construct GlobalVarSupply from IRModule

* Add tests for supply

* Add documentation for NameSupply and GlobalVarSupply

Co-authored-by: Florin-Gabriel Blanaru <fgb@system76-pc.localdomain>
2022-08-11 09:18:46 -07:00
Greg Bonik f5f5a75ae9 [TVMScript] Text underlining in DocPrinter based on Doc's source_paths (#12344)
This adds an ability to print a "diagnostic marker" based on a given ObjectPath. For example, say we are printing a fragment of TIR like
```
for i in T.serial(10):
    a[i] = 5
```
and we would like bring the user's attention to the bound of the loop:
```
for i in T.serial(10):
                  ^^
    a[i] = 5
```
In this case we would give the doc printer an object path that represents this loop bound, i.e. something like `path_to_underline=ObjectPath.root().attr("extent")`

Tracking issue: https://github.com/apache/tvm/issues/11912
2022-08-10 22:43:54 -07:00
driazati 06ac5cde84 [testing] Remove wrapper from @slow (#11566)
This makes it a normal pytest decorator so it doesn't incur test set up / tear down. This also makes the PR body the source of truth for skipping slow tests or not since it can be confusing sourcing it both from the PR and commit message.
2022-08-10 14:35:48 -07:00