main
30 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
 Tianqi Chen
|
c717c5b217 | [IR][Relax][TIRx] Unify Var identity (#20004) | ||
|
Tianqi Chen
|
99869414de |
[TIRX] Remove SizeVar in favor of contextual constraints (#19930)
## Rationale
`SizeVar` encodes nonnegativity in runtime subtype identity, which is
fragile under cloning and remapping. Symbolic integer values should use
one `Var` representation, with nonnegative facts recorded in the
analyzer at the use sites that establish them.
## Changes
- Remove `SizeVar` from the C++, Python, TE, TVMScript, FFI, visitor,
and serialization surfaces, and migrate callers to `Var`.
- Preserve the existing Relax constraint ownership model and use
`MarkGlobalNonNegValue` as the canonical path for global nonnegative
facts.
- Preserve `T.handle()` as the normal opaque-handle form. An optional
dtype constructs a typed pointer, with `T.handle("void")` reserved for
an explicit pointer-to-void.
|
||
|
Tianqi Chen
|
275114b327 |
[REFACTOR][IR] Unify PrimExpr with Expr typed view (#19910)
## Summary - Make `PrimExpr` a typed C++ view over `Expr` values whose `ExprNode::ty` is `PrimType`, instead of using a separate runtime node class as the proof of primitive-ness. - Use the shared `ir::Call` node for Relax, TIRX, and primitive-valued calls, while keeping primitive-only APIs explicit at their semantic boundaries. - Keep Python on the general `Expr` surface for primitive-typed values so `isinstance` behavior does not imply a nominal primitive-expression subclass. ## Design Rationale The main advantage of this change is that common expression nodes such as `Call` can be unified without specializing each one to `PrimType`. A single `ir::Call` can represent a Relax tensor call, a Relax scalar call, or a primitive-valued intrinsic call; the result type stored in `ExprNode::ty` determines whether that particular value can be viewed as `PrimExpr`. This keeps the IR node hierarchy focused on expression structure rather than result-type categories. Nodes that are intrinsically primitive, such as integer and floating-point literals or TIRX primitive operators, still have strongly typed C++ APIs and data structures. General nodes whose result type may vary, such as `Call`, remain general `Expr` nodes and are narrowed to `PrimExpr` only where primitive-only semantics are required. The PR also keeps the compatibility surface practical: C++ primitive-only APIs continue to accept `PrimExpr`, Python exposes a compatibility predicate for checking the primitive typed category, and visitors/printers use one natural `Call` path rather than duplicating Relax and primitive call handling. Missing expression types are represented explicitly with `Type::Missing()` so constructors can leave type inference to later analysis without relying on nullable `Type` values. |
||
|
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. |
||
|
ConvolutedDog
|
e7a7447929 |
[Fix][CI]: remove astral-sh/setup-uv from lint workflow (#19554)
This PR fixes https://github.com/apache/tvm/issues/19552. astral-sh/setup-uv is not on the ASF GitHub Enterprise action allowlist, causing the Lint workflow to fail with "Startup failure" before any pre-commit checks run. See https://github.com/apache/tvm/actions/runs/25743684906 for the failed reason. This PR removes the uv setup and sync steps entirely; pre-commit/action will install and manage pre-commit and all hook dependencies on its own. This PR also corrected previous lint errors. After the fix, the CI lint succeeded: https://github.com/apache/tvm/actions/runs/25775499703/job/75707088129 |
||
|
Shushi Hong
|
3bc61d1fab |
[BugFix][TOPI] Fix get_const_tuple hanging indefinitely when passed a te.Tensor (#19380)
This pr fixes #18765: `topi.get_const_tuple` hangs indefinitely when passed a `te.Tensor` instead of a shape tuple and adds a type check to raise a clear `TypeError` with a helpful message suggesting `get_const_tuple(tensor.shape)` instead |
||
|
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 |
||
|
Tianqi Chen
|
9a8320acbd |
[LINT][PYTHON] Modernize annotations with ruff UP rules (#18830)
This PR enables ruff pyupgrade (UP) rules with py310 target, auto-fixing ~5600 annotation modernizations (PEP 585 generics, PEP 604 unions, deprecated typing imports). Also removes from __future__ import annotations from ir/module.py and rmsnorm.py, bumps requires-python to >=3.10, and removes absolute_import aliases from topi/contrib files. |
||
|
Tianqi Chen
|
33dcea1686 |
[REFACTOR][LINT] Modernize ruff config (#18810)
This PR removes the extra lint violations from the codebase so lint aligns with the latest style |
||
|
Tianqi Chen
|
aa2e609136 |
[LINT] Modernize lint to use pre-commit hooks (#18807)
This PR migrates existing lint to use pre-commit hooks |
||
|
Tianqi Chen
|
c8140643d3 |
[REFACTOR][S-TIR] Move remaining data structures to s_tir (#18743)
This PR moves remaining related data structures to s_tir. - Moves sblock_dependency_info and sblock_scope. - Moves related analyssis. - Hides the data_type_rewriter to private functions. |
||
|
Qingchao Shen
|
eca92bd4f2 |
Fix IR generation conflict in topi.nn.simplify by separating Tensor and PrimExpr handling (#17978)
Update utils.py |
||
|
Tianqi Chen
|
41c9c3b91a |
[REFACTOR][TIR] remove legacy tir::any (#17783)
This PR removes legacy tir::any which was used to represent unknown shape in relay. As we move toward first class symbolic shape, we no longer need the ? shape in the system. |
||
|
Krzysztof Parzyszek
|
71d3262e90 | [TOPI] Use f-strings for string formatting, NFC (#14839) | ||
|
AndrewZhaoLuo
|
95d2e9fa35 |
[ARM] Add dynamic matvec support (#13502)
* [ARM] Add dynamic matvec support * proper imports Co-authored-by: Tristan Konolige <tkonolige@octoml.ai> |
||
|
Junru Shao
|
b16a64d6ed | [MetaSchedule] Refactor ScheduleRule Attributes (#13195) | ||
|
Gavin Uberti
|
e7a72af0c1 |
[microTVM] Improve code reuse in Corstone300 conv2d tests (#13051)
* Add support for out_layout to tensordot schedules and tests * Move shared conv2d test logic into new file * Rework depthwise and grouped convolutions to use common logic * Linting and bugfixes * Fix tests * Fix depthwise and grouped tests * More linting fixes * Address code review comments * Fix unit tests * Address code review comments from Andrew * Fix imports |
||
|
Egor Churaev
|
8d2e887dbb |
[HotFix] Fix python import (#13099)
* [HotFix] Fix python import Tuning doesn't work after #12969. It reports the following error: ``` ImportError: cannot import name 'get_const_float' from partially initialized module 'tvm.topi.utils' (most likely due to a circular import) ``` In this commit I moved import relay to a function which used in a test. And it helps to fix this circular import * Fix lint |
||
|
Gavin Uberti
|
fcbcd156c7 |
[microTVM] Add Cortex-M DSP schedules for optimal conv2d layouts (#12969)
* Rewrite conv2D to tensorize with tensordot * Functional conv2D tensordot implementation * Add stupid hack to work around TVM bug * Unit testing for conv2d schedule * Connect new implementations to Arm strategy * Separate into new tensordot conv2d schedule * Separate testing infrastructure * Prototype depthwise implementation * Unit testing for depthwise_conv2d * Linting and documentation * Enforce SIMD alignment in strategy * Prevent black from butchering our formatting * Address code review comments * Fix alignment strategy bug * Fix linting * Remove unconventional offset behavior * Replace math.prod function to support Python 3.7 * Fix CI tests |
||
|
Masahiro Masuda
|
a34731b7fc |
[ROCM] DP4A intrinsic support for TE/TIR (#11009)
* [ROCM] Support dp4a on AMDGPU by sdot4 intrinsic commit 0225f2bfe3f413cd4764c2dba6c922af2520146b Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 14 08:56:10 2022 +0900 share op strategy between cuda and rocm commit 762c7e8611c9ec3cca3321428e2362c81fe89b9b Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 14 08:28:34 2022 +0900 fixed rocm batch_matmul strategy for mixed i8i8i32 commit ce53e8d141f7f901303ec6a91674337cbf2b2384 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 14 06:17:30 2022 +0900 add rocm sdot4 TIR intrin commit f4562b991f9180b61be7339b2890de1584656c10 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 14 06:03:44 2022 +0900 rocm sdot4 works commit 6cc62805f82dd884a18a1c4c0e9bae5866e00da0 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 14 05:32:07 2022 +0900 more wip commit 0602f4a3157d4cb5a3f280a3a3c514bb6535aac8 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 14 03:47:37 2022 +0900 Squashed commit of the following: commit 65b8bcf955f44540d6a52c8416e60f3047c8366c Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 13 20:36:49 2022 +0900 [WIP] adding DP4A support to rocm commit 4f8f308ab6bb85ef3bdcc2b8e846c2eea15f2167 Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 13 14:03:25 2022 +0900 Squashed commit of the following: commit 1711be38a17e3b6171350009f1da05824cd0b340 Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 13 13:11:40 2022 +0900 fixed condition for real commit 8a48fb5262e80e318cd81d5ff51bf95fd5eb576e Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 13 09:57:42 2022 +0900 Revert "Skip applying sch_rule when both ann and sch_rule are defined" This reverts commit 4915c6a5a91ff87038e71f8aff9f31db684b4a95. commit daea033d2cb06388ef27ddadb80fc5bce72181d2 Author: Masahiro Masuda <masahi129@gmail.com> Date: Mon Apr 11 09:31:05 2022 +0900 [Metaschedule] Support rocm and spirv commit eb0cae2c779808cced074d189e8f487bf46ea89f Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 13 07:25:04 2022 +0900 dp4a works commit 4915c6a5a91ff87038e71f8aff9f31db684b4a95 Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 13 06:13:45 2022 +0900 Skip applying sch_rule when both ann and sch_rule are defined commit 7b3d71c6b21a9c5de9ef2b89d0a7db2800a5f3a2 Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 13 04:40:31 2022 +0900 fixed intrin description commit 7666cd7a5b0ce182791662673fbe45944c84d0ae Author: Masahiro Masuda <masahi129@gmail.com> Date: Tue Apr 12 19:59:47 2022 +0900 add DP4A intrin commit 7086bdb75546a2680d12dc8f80c040cea23f729a Author: Masahiro Masuda <masahi129@gmail.com> Date: Tue Apr 12 19:03:44 2022 +0900 works commit db343974bfae86e51078e40e6170022a782d8e0a Author: Masahiro Masuda <masahi129@gmail.com> Date: Tue Apr 12 12:49:52 2022 +0900 more hack to tensorize loop mapping to make resnet50 e2e work commit 2409674a7884a60beb50d7aa3345c4b907b8cd13 Author: Masahiro Masuda <masahi129@gmail.com> Date: Mon Apr 11 13:40:59 2022 +0900 wip support pad + qnn.conv2d folding commit 613cb7ec33b6df41f1ebe0f0a0ac8eca7c73cff1 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Apr 10 12:04:08 2022 +0900 hack to tensorize loop mapping to make conv2d work commit 9e4f9df6a409396a8a4a20d967c4f51accf5d210 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Apr 10 11:34:13 2022 +0900 wrap tensorize with try/catch commit d4b496d858da0ae43063d47cb03a28b803d0269f Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Apr 10 11:33:39 2022 +0900 revert change in task_scheduler.cc commit 476129be7b286f5d109402280aea585e89f6dc1d Author: Masahiro Masuda <masahi129@gmail.com> Date: Sat Apr 9 05:54:10 2022 +0900 try / catch in ThreadedApply commit d8226ff26f25eba17d4000f25131822874bdc2cc Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Apr 8 17:17:59 2022 +0900 filter out invalid candidate commit 2632899a2759885d338e25f2a25ba0b2c555f0c3 Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Apr 8 10:09:48 2022 +0900 try graceful exit in parallel_for_dynamic commit 9d6741c3dd29c4dde861aa1d3b2ca85f560f5ac6 Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Apr 8 09:35:51 2022 +0900 [QNN] Fix broadcast for invalid axis commit 6ccde0959343ce4246ef99505b4f54de469a1a5c Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 20:51:15 2022 +0900 refactor rewrite_tensorize commit 2ce206699f10b03b9611c4683018f7e0c70c7eb5 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 20:48:17 2022 +0900 allow missing schedule_rule in post order apply commit 3a69353a29abfc454e28d4e530d22a3e2043712e Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 19:42:48 2022 +0900 refactor rewrite_tensorize commit 43e0b2f7f98299679807aaf1ffb13cce2b5f5ce3 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 18:25:14 2022 +0900 rewrite_vnni -> rewrite_tensorize commit 823797e2627a9bfa812b72019468569ee79eb4c6 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 18:12:12 2022 +0900 VNNI -> WithIntrin commit 4284a47e5933aa89c1c3362b15ad53b14782fc81 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 17:45:41 2022 +0900 introduce TileForIntrin commit b87ef32e30e1e71b3f39789f7289976a8cba4ab4 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 17:34:04 2022 +0900 move TilingwithTensorIntrin to auto_tensorize.cc commit 2fc118b3726586ba13f7de950beaa299b83a0af3 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 17:28:45 2022 +0900 clean up headers commit d8b2aa325c91b524bec22dc1ec2fc52c9f060fce Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 17:09:32 2022 +0900 clean up using namespace commit eb05d25e2b71f4a1232a8796d1413011ec7629d3 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 17:03:05 2022 +0900 refactored init commit 5e6b0a08d447c0470c2c8a993e4bd62673e34fe3 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 16:57:14 2022 +0900 compiled commit 2b8c430e2fec7ceb285eed7bc7aa73bb9a74a997 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 12:51:55 2022 +0900 wip MultiLevelTiling refactor commit 7c21a9fea0511c88bd82f49f799b5198252df40a Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:58:33 2022 +0900 function doc string not supported by tvmscript commit 40f9742bc9c3aa11e8c2c0551d1827ad47fc0f39 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:56:45 2022 +0900 update vnni intrin name commit 4814f825a5315efd2a3da8c36d2ce6b5df5447cd Merge: e0c5eb84b 07bbb38f7 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:44:47 2022 +0900 Merge branch 'tir-tensor-intrin' into auto-tensorize-vnni commit 07bbb38f7fb52db4a2ecde3d5c87cf4d5cd000a1 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:24:56 2022 +0900 more lint fix commit 15e60b42362cc64b1428b219c8eada414d1b8372 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:16:08 2022 +0900 black commit 7a757fe53758e06418ea1367b348b47c8cd2dcf9 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:12:54 2022 +0900 pylint commit 9a3e508b6f4529158e703b4617f2ddaa351a89eb Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 10:58:52 2022 +0900 simplify import commit d8e43ecf1c0a79a2c195ff31e1e699a447a11335 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 10:52:50 2022 +0900 use vectorlow/high in arm intrin commit 625cd2774ec455307646b0c26bb3971d89613d1e Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 10:34:57 2022 +0900 fixed offset factor commit 69e72b6b612588e670937e003435afa647030ceb Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 10:12:02 2022 +0900 Add ARM intrin commit 1351fdea6b22f231a290a6c28e06732c9cf993cf Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 08:27:27 2022 +0900 use buffer syntax sugar commit 0ced85fd097ed48aad8714912718d8735791e1fb Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 08:17:43 2022 +0900 rename vnni.py to x86.py commit 38a5aca87ec438446593a3af17760339211f5ad9 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 07:24:44 2022 +0900 add VNNI unittest commit 88b763ec48c20cf68db8bc3bae3fa3ae78996ee8 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 07:10:06 2022 +0900 refactored existing test using VNNI intrin commit 711a0076d9be2b9aa80ada67e1edda5ba1fdf1fd Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 07:04:58 2022 +0900 [TIR] Add VNNI dot product intrinsic for TIR commit e0c5eb84bf6a0ad2ba0cddc4bdf22a799dc4b8a0 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:42:26 2022 +0900 merge fix commit b171748139e53f0cf75ff4b6fde436f9d8a5fe91 Merge: 71fe3bdf0 82e152a3c Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:33:59 2022 +0900 Merge branch 'tir-tensor-intrin' into auto-tensorize-vnni commit 71fe3bdf02ae10ddbe090a4fd1020f545a05bb41 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 06:57:38 2022 +0900 move tensor intrin under tir commit 0c51badef45af2a1025ab42fe38d1b3f07ab493e Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 06:12:39 2022 +0900 remove log commit fed910e03eb94c169d4a160b8f3cad406d04c6aa Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 06:11:22 2022 +0900 more revert commit 7150aff9fba167d88dbfb40d48727de8a144b9c0 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 06:10:44 2022 +0900 revert stmt_functor change commit 155107b98b09c5e5cc7f19afbd327b0557a02843 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 06:10:09 2022 +0900 refactored RewriteVNNI a bit commit ca15255e3a882b89b05bb83079640c929fb63096 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 05:41:13 2022 +0900 add RewriteVNNI commit dc9f71d5e3122b50fa8ae6a4462f959f13870b05 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 05:38:56 2022 +0900 vectorized init loop commit fcc31ee20ddfafd47f566bf98ff40a9f684d12eb Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 04:55:36 2022 +0900 tensorize worked commit 2b534377a45b9ab84bf35c3d7c03ecae7616d17f Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 6 19:11:05 2022 +0900 TilingwithTensorIntrin works commit 86baa31e773fc864f77dc113bc9a93b79f3fc652 Author: Masahiro Masuda <masahi129@gmail.com> Date: Wed Apr 6 08:58:27 2022 +0900 Ported auto-tensorization code commit 82e152a3c91144041ade783116a50565ebb48b89 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:24:56 2022 +0900 more lint fix commit 88d9bdd3b21302bc2dd068a990df15c375a1a8ef Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:16:08 2022 +0900 black commit 31fe7eb8075445161d804d170772eac8e90d3425 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 11:12:54 2022 +0900 pylint commit 7876754effc40ad089349534dacd75df19d38fc4 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 10:58:52 2022 +0900 simplify import commit 56f2e9a85069426021e2872eb1da95bf134ac7e0 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 10:52:50 2022 +0900 use vectorlow/high in arm intrin commit 995cc8d6fcec70a3fadcfb1c6fee7b9f0b5a0951 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 10:34:57 2022 +0900 fixed offset factor commit 86bbd4955b34257d68d957cb4a2536aea3ef9bac Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 10:12:02 2022 +0900 Add ARM intrin commit 120fd96e80307b4301ee3fc93e6793e0b40485f0 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 08:27:27 2022 +0900 use buffer syntax sugar commit 0f0682d00c3961afd1f492ae55f180c5b5502767 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 08:17:43 2022 +0900 rename vnni.py to x86.py commit f88c31ead1fa6db4bfd2c88eeaf5f665e4c6dddb Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 07:24:44 2022 +0900 add VNNI unittest commit 6cc80094adac398762924b0b31a4c741417ba9dc Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 07:10:06 2022 +0900 refactored existing test using VNNI intrin commit 11a29c704cdaad96aeeca39c9c753ef006d27a50 Author: Masahiro Masuda <masahi129@gmail.com> Date: Thu Apr 7 07:04:58 2022 +0900 [TIR] Add VNNI dot product intrinsic for TIR * cleanup * black * update dot prod intrin * add mattr kind * conv2d topi test working * add dense and bmm test * add conv2d relay test * add tir intrin test * pylint |
||
|
Eric Lunderberg
|
174d09ee2c |
[TE][TIR] Implement layout transformations, non-flat memory buffers (#9727)
* [TIR] Added BufferLoadNode::LegalizeDtype
When modifying a BufferLoad object, the return dtype must also be
updated. This exposes the legalization function, so that passes that
use `BufferLoad::CopyOnWrite` to modify the buffer/indices don't need
to repeat the logic to update the dtype returned.
* Replacing Store/Load in Stmt/Expr Visitor/Mutator
* Removing Store/Load from optimization passes
- UpdatePointerStorageScope
- UnrollLoop
- ThreadSync
- LinearAccessPatternFinder
- StoragePlanRewriter
- VectorTypeRewriter
- VectorTypeAccessChecker
- NarrowDataType
- IRConvertSSA
- CompactBufferRegion
* Removing Store/Load from examples
- ConvertAddToSubtract
* Replacing Store/Load in StorageFlatten
Now, outputs BufferLoad/BufferStore with a flattened buffer object.
temp commit, replacing Store/Load, BufferBindUnwrapper
temp commit, replacing Store/Load, StorageFlattener
* Replacing Store/Load in utility passes.
- StmtSimplifier
- IRSubstitute
- BaseInliner
- FeatureVisitor
* Replacing Store/Load in analysis functions
- StorageAccessVisitor
- VarTouchedAnalysis
- MemoryAccessVerifier
- InplaceOpVerifier
- GPUCodeVerifier
- VarTouchVisitor
- LCADetector
- BlockReadWriteDetector
- InstrumentBoundCheckers
* Replacing Store/Load in lowering/legalization passes.
- MakeCrossThreadReduction
- CacheReadRewriter/CacheWriteRewriter
- InjectVirtualThread
- InjectDoubleBuffer
- InjectCopyIntrin
- LowerWarpMemory
- LowerThreadAllreduce
- LowerThreadAllreduce
- LowerCustomDatatypes
- LowerTVMBuiltin
- CoProcSync
- MergeDynamicSharedMemAllocations
- VectorizeLoop
- BF16Legalize
* Replacing Load/Store in codegens.
- Device code generators
- CodegenC
- CodegenLLVM
- CodeGenOpenCL
- Utilities used during codegen
- ArgBinder
- MakePackedAPI
- ReturnRewriter
- SplitHostDevice
- Execution environments
- CodeGenStackVM
- CodeGenHybrid
- AOTExecutorCodegen
* [UnitTest] Add unit tests to test physical layout remapping.
* Updated tvm::address_of() to hold BufferLoad instead of Load.
* [TIR] Added IndexMap class.
Holds a set of variables representing the input indices and
expressions in terms of those input indices.
TODO:
- Add validation, the index mapping should be invertible.
- Add helper function, apply mapping to a set of indices.
- Add helper function, apply mapping to bounds of input indices.
* Updated Buffer::vstore/vload to return BufferLoad/BufferStore objects.
StorageFlatten/FlattenBuffer passes updated to modify the
buffer/indices directly, rather than using vload/vstore.
- Primary purpose of vstore/vload is to allow IR written in python to
define vectorized load/store. This usage is maintained by returning
a BufferLoad/BufferStore node whose index is a Ramp.
- Previously, vstore/vload was also used to compute the 1-d physical
index of a location within a N-d tensor. This usage will no longer
be allowed, as it would not allow layout transformations to be
performed after a schedule definition, but any uses of the buffer
are flattened.
* [TE] Added Stage::transform_layout to the C++ TE implementation.
Adds an `Array<IndexMap>` in the stage to define the transformations
to be applied on the tensor's layout. As of this commit, this mapping
isn't propagated into the TIR graph yet.
* Replace Store/Load with BufferStore/BufferLoad in ir_builder
* [TE] Added Stage.transform_layout to the Python TE interface.
Allows users to specify `s[A].transform_layout(mapping)`, and
propagate into the TE definitions.
* Added pre_flattened_shape/pre_flattened_stride fields to Buffer.
The shape and stride checks performed in ArgBinder::BindDLTensor
(called from MakePackedAPI) require the tensor shape/strides prior to
index flattening. Therefore, though it is no longer used by the
low-level code generators, we must maintain that information for use
in MakePackedAPI.
* [UnitTest] Test N-d indices exposed to low-level codegen
When using te.AXIS_SEPARATOR in the call to .transform_layout, this
should define groups of axes, each of which is flattened to a single
axis, then exposed to the low-level codegen.
* [TIR] Added PrimFunc attribute "layout_transform_map", filled from TE.
Propagated the TE definition of the physical layout into the TIR
graph.
* Added pre_flattened_type.
If a boolean tensor is backed by an int8 buffer, the check on the
argument buffer's type should be against the boolean type.
When rebasing this PR, should be placed after the addition of
pre_flatten_shape/pre_flatten_strides.
* [UnitTest] Added tests for loop iteration order.
After transformation, the iteration order should follow the new
transformed axes. In addition, the loop iteration variables should be
exposed through the TE interface for further manipulation.
* [TIR] Added BufferNode::axis_separators
- Add axis_separators to represent divisions between groups
of tensor axes, where each group is flattened into a single
output axis, to be exposed to the low-level code generators.
- Expose axis_separators to the python interface.
- Update existing C++ calls to the Buffer() constructor.
* [TIR] Added ApplyLayoutTransforms as part of StorageFlatten.
For any buffers that have layout transforms defined in the
"layout_transform_map" attribute of a PrimFunc, rewrite access into
the buffer such that they use the updated ordering.
* Update usage of ir_builder where necessary.
* [TE] Implement te::Transform
Similar to Fuse and Split, this represents a modification to the
existing loop iterations.
* [TE] Added Stage::set_axis_separators.
In C++, this is implemented as an `Array<IntImm>`, specifying
pre-flatteneing axes after which a new post-flattening should be
started. The python interface uses a sentinel value
`te.AXIS_SEPARATOR` in the call to `transform_layout`, which is then
used to define the array of axis separators.
* [TIR] Expose tir.transform.ApplyLayoutTransforms for testing
* [TE] Rewrite loop iteration order
After .transform_layout, rewrite leaf_iter_vars to follow the updated
order. Use the te::Transform iter_var relationship to track use of
the transformed variable.
* [TE] Fill BufferNode::axis_separators from StageNode
During ScheduleOps and SchedulePostprocToPrimfunc, the axis separators
defined in the stage must be passed through to the TIR BufferNode.
* [TE] Return transformed iteration variables
* Moved Buffer's pre-flatten information to PrimFunc.
Since the pre-flatten information is only used for validating user
inputs, it makes much more sense to store it alongside the buffer_map.
* Updated ethos-u C++ unit tests to remove use of Load/Store.
* Bugfix, layout transformation.
Error occured during conversion from TE to IRModule, when layout
transforms were applied to a reader of a `cache_read`.
* In test directory, replacing all instances of T.load.
* Return buffer object from tvm.tir.script.scope_handler.Allocate
Now that the load/store require buffer objects, allocation should also
return a buffer object to be used.
* Added .astype to tvm.script.tir.node.BufferSlice
Since `buf[i]` returns a `BufferSlice`, this lets the TIR examples
that use `buf[i].astype('out_dtype')` continue functioning.
* Replacing all T.store TIR calls.
* Added LOG(FATAL) in constructor of Store/Load nodes.
* Updated tvmscript parser to report error for Store/Load nodes.
* [TVMScript] Added T.preflattened_buffer stmt
Used to specify `PrimFunc::preflattened_buffer_map`. Takes an argument
of the postflattened buffer, so that it will work for both simple
declarations and `T.match_buffer` statements without needing to
introduce a param handle. All other arguments are identical to
`T.match_buffer.`
* [TVMScript] Updated TVMscript for BufferLoad/BufferStore
- Use `T.preflattened_buffer` calls in TVMScript to represent
`PrimFunc::preflattened_buffer_map`.
- Remove `T.buffer_decl` for return value of `T.allocate`, now that
`T.allocate` returns a buffer.
- For buffer access as a different type, make a `T.buffer_decl` for
those accesses.
* Updated test_tvmscript_roundtrip.py for BufferLoad/BufferStore.
* Updated TIR reference in USMP pool allocation unit tests.
Using let var handles as the data pointer in buffers, rather than just
as `T.load`/`T.store` arguments, requires annotation as
`T.Ptr[T.primtype]`, rather than as `T.handle`.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* fixup! Replacing all T.store TIR calls.
* fixup! Replacing all T.store TIR calls.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* fixup! In test directory, replacing all instances of T.load.
* tir.ComputeInline, correct variable count.
Previously, this metaschedule primitive relied on `tir::UndefinedVars`
ignoring the data pointer of BufferLoad/BufferStore nodes. When
`tir::UndefinedVars` was updated to visit the data pointer, similar to
the previous behavior when visiting Load/Store nodes, this caused the
count of undefined variables to be unexpectedly high.
* fixup! Replacing all T.store TIR calls.
* fixup! Updated Buffer::vstore/vload to return BufferLoad/BufferStore objects.
* fixup! In test directory, replacing all instances of T.load.
* fixup! In test directory, replacing all instances of T.load.
* fixup! Replacing all T.store TIR calls.
* Expose Buffer index flattening function to Python.
* Updated test_tir_buffer.py offset tests.
Replacing calls to `Buffer.vload` with `Buffer.offset_of`, when
testing the index calculations.
* fixup! Replacing all T.store TIR calls.
* fixup! Replacing all T.store TIR calls.
* fixup! Updated Buffer::vstore/vload to return BufferLoad/BufferStore objects.
* fixup! Replacing Store/Load in lowering/legalization passes.
* fixup! Replacing all T.store TIR calls.
* fixup! Updated ethos-u C++ unit tests to remove use of Load/Store.
* fixup! Replacing Store/Load in lowering/legalization passes.
Fix linting for inject_double_buffer.cc
* fixup! Updated ethos-u C++ unit tests to remove use of Load/Store.
* fixup! Added .astype to tvm.script.tir.node.BufferSlice
* fixup! In test directory, replacing all instances of T.load.
* fixup! Replacing all T.store TIR calls.
* fixup! Replacing all T.store TIR calls.
* fixup! In test directory, replacing all instances of T.load.
* fixup! Replacing all T.store TIR calls.
* fixup! Replacing Store/Load in lowering/legalization passes.
* [UnitTests] Added T.preflattened_buffer in expected result
* fixup! In test directory, replacing all instances of T.load.
* [UnitTests] Bound checker update, compare against N-d buffer bounds.
* Fixup, bound checker vectorize test.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* [UnitTest] Fixed breakage in InjectRollingBuffer test.
Needed a bit more re-writing than usual, because the test was
explicitly calling lowering passes, then calling `tvm.build`. Fixed
by using the standard lowering flow, with preprocessing steps
inserting with `tir.add_lower_pass`.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* [UnitTest] Fixed breakage in flatten buffer unit tests.
- Updated pass to allow BufferStore/BufferLoad nodes to be visited
before the block's alloc buffer.
- Added `T.preflattened_buffer` annotations.
* fixup! Return buffer object from tvm.tir.script.scope_handler.Allocate
* [UnitTests] Fixed breakage in test_tir_buffer.py
- Updated vload test for new behavior.
- Added test for offset_of, testing behavior no longer in vload.
- Added null check for buffer visitor.
* fixup! Replacing Load/Store in codegens.
* [UnitTest] ComputeInline, opaque access test updates
* [UnitTest] Fixup, allow unit test to use `ib.pointer()[0]`.
* fixup! Replacing Load/Store in codegens.
The updated CodegenLLVM should use the BufferStore/BufferLoad
convention of indexing by `sizeof(dtype)`, rather than
`sizeof(dtype.element_of())`.
* fixup! Replacing Store/Load in lowering/legalization passes.
BF16Legalize should also update the preflattened_buffer_map, since it
is overwriting the `BufferNode::data` stored in the buffer_map.
* fixup! Replacing all T.store TIR calls.
* Fixed failing codegen c host unit tests.
- Generated functions were making `uint8_t*` parameter arguments for
array handle for return value, rather than the earlier `void*`.
- New parameter type was due to using
`PointerType(PrimType(DataType::UInt(8)))` as the type annotation, to
be usable as `BufferNode::data`.
- Changing to `PointerType(PrimType(DataType::Void()))` still allows
usage as buffer, more appropriately expresses semantics.
- Updated C codegens to allow `void*` types to be generated from
variables with type annotation, in addition to the previous behavior
of `DataType::Handle()` variables without type annotation.
* Fixup, StorageFlatten when applied to post-StorageRewrite functions.
Identified in a test that applied `tvm.lower`, then `tvm.build` on the
result. If the result of an allocate node is used as the backing
buffer for multiple buffers, such as the output of the StorageRewrite
pass, then StorageFlatten would erroneously think that the second
occurrence was an usage without earlier definition.
* fixup, StorageFlatten
When flattening a boolean buffer, the backing buffer should have type
int8, not the preflattened buffer.
* Bugfix, correctly represent void* in LLVM IR.
* Update, replace tir.Load with tir.BufferLoad
* Added TVMScript error check for matching buffer/index dimensionality
Needed for tests/python/unittest/test_tvmscript_error_report.py::test_high_dim_store
* Bugfix, correct return type when lowering custom datatype.
* Bugfix, removed unused primfunc from test_tvmscript_complete.py
* Updated test_meta_schedule_postproc_verify_gpu_code.py TIR
Replaced Load/Store with BufferLoad/BufferStore.
* Allowed ramp nodes with buffer use analysis.
* Updated tests in test_meta_schedule_postproc_verify_gpu_code.py
Needed dummy writes to prevent buffer resizing, in order to trigger
the verification failure due to memory limits.
* Updated TIR examples to be compatible with buffer dimension check.
* Corrected section header in docstring.
* Corrected indices size check in CogeGenC.
* Fixed breakage in LowerThreadAllreduce.
Since the AllocateNode is rewritten, any buffers that refer to those
variables must also be rewritten.
* [UnitTests] Replaced Store/Load in CUDA codegen tests.
* Resolved breakage in C-based codegen for vectorized store/load.
Needed to update to new convention of using the buffer's element type
as the stride.
* Bugfix, incorrect LCA for buffer access in root scope.
This had been present before the BufferLoad/BufferStore changes, but
hadn't triggered on tests using Load/Store nodes.
* Added docstrings for TransformNode member variables.
* Added TODO for future removal of preflattened_buffer_map.
* Fixup, transform layout + cache write tests.
The correct sequence is to first apply any caching as needed, then to
apply layout transformations, and finally to apply thread binds for
the computation step.
* Bugfix, correct element type for scalarized access.
* Bugfix, cuda buffer indexing when declared as different type.
* Cuda codegen, update reference.
* Bugfix, lower allreduce
Loads of the output of the reduction should be replaced for all
buffers sharing a buffer pointer, not just for the buffer object
itself.
* Removed obsolete comment.
* Changed PrimFunc constructor preflattened_buffer_map to Optional
* Removed flatten_buffer argument from T.match_buffer.
* Correct call to VarUseDefAnalysis::VisitBuffer
* Reverted unintentional testing change, lanes=2.
* Updated lower_cross_thread_reduction to use buffer in allreduce
* Updated transform_layout test to disable CSE
* Updated CSE unit tests to use BufferStore
* Replaced Store/Load for vta.transform and unit tests.
* Updated unit tests for lower_cross_thread_reduction.
* Updated arange to use scalar tensors.
The start/stop/step tensors are declared as 0-d scalar tensors, but
were accessed as 1-d tensors.
* Fix breakage in ethosu constant encoding.
Buffers generated by "ethosu_copy" should have their buffer objects
rewritten, but shouldn't have their size updated in ethosu-specific
Call nodes.
* Fix breakage in ethosu call argument checks.
Need to pull out indices from BufferLoad holders, not Load.
* Resolve breakage from mismatched shape/index dimensions
* Split out encoded parameters from preflattened buffer map.
* Updated buffer shape/index dimensions to match in more ethosu tests
* Fixed lint error
* Removed debug code
* Moved arith::Analyzer local variable to class member
* Fixed SSA conversion of allocations.
Can occur if allocation is inside an unrolled loop. Added unit test
to catch this failure mode.
* Ethos-u index/buffer dimension updates.
* Updated ethosu passes to handle buffer load/store.
* Resolved bug in tvmscript printing of duplicate buffers.
* Fix breakage in ethos-u test_assign_addresses, encode constants
* Apply same changes to T.allocate_const as to T.allocate
Return a buffer when used in TVMScript, allow for aliasing buffers.
* Fix lint errors.
* Further updates for ethos-u tests.
* Updated ethos.u buffer sizes in test.
* Updated tir.BindParams to use BufferLoad instead of Load.
* Updated topi.cuda.scan implementation to follow buffer dimensions.
* Resolved breakage when flattening AllocateConst nodes.
* Resolved breakages from latest merge with main.
* Corrected error in merge.
* Use empty indices for rank-0 tensor.
* Added ir_builder workaround for 1-d indexing.
* Consistent buffer access type in LLVM codegen, to match C codegen
* StorageRewrite, update indices of modified buffers.
* Dynamic relay nodes, access 0-d tensors with 0-d indices.
* BFloat16 legalization, update buffer type.
* Updated meshgrid to use 0-d index for 0-d buffer.
* Corrected boolean handling in Allocate nodes.
* Added workaround to unpack 1-d Tensor indices into N-d buffer indices.
* Resolved a few more failures in relay tests on cuda.
* Resolve linting
* CI bump
* Updated renormalize_split_pattern tests to use BufferLoad/BufferStore
* Fixed cuda codegen checks for BufferStore/Ramp.
* Simplify indices further, needed to avoid cuda register limit.
* fixed dyn onehot shape func accessing 1d buffer with ()
* Fixed codegen indexing for int4 scalar types.
* Temporary workaround for incorrect constant folding.
Need to further investigate vectorized LLVM constants
* s/find_allocate_usage/FindAllocateUsage/g
* Added buffer type consistency TODO.
* Improved comment on address_of Op.
* Rename LegalizeDtype to LegalizeDType, made private.
* fix format and lint errors
* Disable vectorization of AllocateConst buffer in StorageRewrite.
* Pass buffer_map through to the PrimFunc in cmsisnn
* try disabling problematic winograd test case
* try different way of buffer mapping in storage_rewrite
* Removed unnecessary ramp node in ir_builder.
* Updated LLVM codegen for buffer indexing.
TVM data arrays are always densely packed. If the LLVM type
corresponding to a vectorized TVM datatype contains padding for
alignment, the array location should be computed based on the
primitive element type.
Co-authored-by: Masahiro Masuda <masahi129@gmail.com>
Co-authored-by: adstraw <astraw@octoml.ai>
|
||
|
Junru Shao
|
865f34ea5f |
[MetaSchedule] Enable AutoTVM-style template-based search space (#10461)
* [MetaSchedule] Enable AutoTVM-style template-based search space * Fix lint * suppress mypy |
||
|
Matthew Brookhart
|
d3fc562a6f |
[Relay][TOPI] Resize 1D (#8346)
* rename resize to resize2d * refactor resize_2d * Add resize1d op, normalize attribute names across ops * normalize resize3d to match the API of 1D and 2D * fix lint * fix relay tests from API change * refactor topi tests, docs * fix method naming in framework frontends fix more frontend issues * refactor resize tests to reuse components, add more coordinate tranform modes to tests * add cubic resize reference kernel and tests, add relay tests for resize1d * fix pylint * fix test typo |
||
|
Mehrdad Hessar
|
a1cd6d51b8 | fix py files (#8194) | ||
|
Masahiro Hiramori
|
84359a97bd |
Fixed minor misspelling (#7499)
Co-authored-by: mshr-h <mshr-h@users.noreply.github.com> |
||
|
masahi
|
1e0d3569b9 |
[Relay, TOPI] Add numpy style cumsum op (#7334)
* Add cumsum relay/topi op * relay tests working * add torch frontend converter * fix for importing detr * fix bad merge * begin cuda cumsum * support non innermost axis * support rank higher than 3 * making binop parameter * fix overflow issue in thrust scan * generic binop parameter working * relay test working * fixed for bool input * remove pytorch change * fix pylint * doc update * Update python/tvm/topi/cumsum.py Co-authored-by: Tristan Konolige <tristan.konolige@gmail.com> * Update tests/python/relay/test_op_level3.py Co-authored-by: Tristan Konolige <tristan.konolige@gmail.com> * add example outputs * add supported input and output dtype in thrust log * adding more loop var names * fix cpplint * fix missing check for the cuda target in nms thrust sort * parallelize cpu cumsum * making binop argument tir function * update doc for binop * doc update Co-authored-by: Tristan Konolige <tristan.konolige@gmail.com> |
||
|
masahi
|
62f251bb34 |
[TOPI] Make cumsum IR reusable, add thrust scan (#7303)
* import changes from scan branch commit cf0d4fdf3bf8fa6e1d6abf631042de28176923c3 Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 25 10:12:01 2020 +0900 get valid count test working commit eb142d3ee9bb16ddf8d37fdec10c1bcda209deaa Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 25 07:22:00 2020 +0900 integrate new cumsum change commit f89684d73dad1f863b4fd291e8804b5c24eae94f Author: Masahiro Masuda <masahi129@gmail.com> Date: Fri Dec 25 06:56:46 2020 +0900 remove ceil_div from nms commit a2ad4dea87d9a637745fb0a40ff9bbdde286194a Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 20 20:36:34 2020 +0900 add api for returning reduction from ex scan output commit b7f4ef7006b722e365533bec53b1f104aa056da2 Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 20 19:49:07 2020 +0900 move ceil_div to utils commit a9a57e34317b1f254165c3a88e465e33c7fda01b Author: Masahiro Masuda <masahi129@gmail.com> Date: Sun Dec 20 19:38:15 2020 +0900 rename prefix_scan.py to scan.py commit 03ed43ff550a435a28740ce1fa62cea71b90cf2c Author: Masahiro Masuda <masahi129@gmail.com> Date: Sat Dec 19 06:12:55 2020 +0900 surpress cpplint commit abceac980d8dfd94072acc228108d1fcd94a214c Author: masa <masa@pop-os.localdomain> Date: Fri Dec 18 20:36:24 2020 +0900 support more data type commit 3e7d1f81821a1e221cbb1322ef5b23f273f51c42 Author: masa <masa@pop-os.localdomain> Date: Fri Dec 18 20:09:51 2020 +0900 1d thrust scan working commit ac13b407e21a83ca57240cad205c32a5d000f999 Author: masa <masa@pop-os.localdomain> Date: Fri Dec 18 19:49:25 2020 +0900 adding thrust scan support commit 65634e86c33786541485dc6461a96da833332297 Author: masa <masa@pop-os.localdomain> Date: Fri Dec 18 19:01:11 2020 +0900 add thrust scan python stub commit 9876c901ee8b406bc9d75ba91c4734d55f85811b Author: masa <masa@pop-os.localdomain> Date: Fri Dec 18 20:55:14 2020 +0900 introduce prefix_scan.py and move scan ir in nms.py commit 667bdd3b135a03b53937fdb664915e07f1365ee1 Author: masa <masa@pop-os.localdomain> Date: Fri Dec 18 15:06:18 2020 +0900 make the scan loop exclusive commit 480787bc072bfc59dcc279038c772f8ad2ec03e9 Author: mbrookhart <mbrookhart@octoml.ai> Date: Thu Dec 17 10:01:11 2020 -0700 Parallelize cumsum in get_valid_counts * fix for 1d scan * rename * cast to out dtype * do not run return reduction for inclusive scan * remove another ceil_div definition * adding scan test * add scheduling for scan op, fixed scan 1d test * pylint fix * add doc string * add more thrust scan test * add dynamic get valid count test, including empty size tensor * fix hard coded gpu targets for cpu only env * try retunring early if scan_size is 0 * another change for empty tensor and thrust path Co-authored-by: masa <masa@pop-os.localdomain> |
||
|
Lianmin Zheng
|
9a9ec1aea3 |
[AutoScheduler] Accelerate feature extraction for winograd (#6981)
* [AutoScheduler] Accelerate feature extraction for winograd * fix an overflow in feature.cc * address comments * address comments * Update include/tvm/te/schedule.h Co-authored-by: Cody Yu <comaniac0422@gmail.com> * Use a smaller min_repeat_ms * Use a smaller min_repeat_ms Co-authored-by: Cody Yu <comaniac0422@gmail.com> |
||
|
Matthew Brookhart
|
a7b22ab9ba |
Dynamic gpu tests, add dynamic strided slice to topi (#6870)
* enable GPU tests for dynamic ops * strided-slice can't do 0-sized output tensors, remove test * move dynamic strided slice into topi * add python interface to topi dynamic strided slice add python interface, tests * autoformat * fix bad copy/paste * fix doc string * disable topk on gpu for now, remove invalid slice test |
||
|
Tianqi Chen
|
c8064b3ca6 | [REFACTOR] Remainings of util => utils (#6778) |