This PR includes the distributed measurement of tuning candidates using builder and async runner, as well as some auxiliary functions. It enables multiple builders and multiple runners with a tracker connecting in between. The hierarchy of files in the database can be further compacted to make the database more concise.
Previously only static shape computation is allowed in our tuning
system. However, one special case is overlooked: the reduction iter vars
could still have dynamic iteration domains which depend on other data
parallel vars. This PR rules out this case by carefully checking all the
loop extents during task extraction.
Related issue: https://github.com/apache/tvm/issues/11746.
* emptycommit 2nd try
* add operator and test
* example output
* lint with black
* register param index
* remove assert as it is a warning in torch
* fix algo bug
Co-authored-by: yuanfz <42092999+FZYUAN-1@users.noreply.github.com>
Minor modification to scoped timers to cover 99% of all the time cost during MS tuning. Allow `ApplyHistoryBest` and `TaskExtraction` time to be counted during tune_relay.
Follow-up of #11622, per discussion with @Kathryn-cat
- [x] Allow using a string `"default"` in `TuneContext` to quickly specify a set of target-specific
rules
- [x] Enhance detection of `ScheduleFn` in `TuneContext` to make it easier for users to quickly try
out template-driven scheduling on TIR.
Next PR:
- Add `TuneContext.tune` to allow directly tuning without task scheduler.
Co-Authored-By: Kathryn (Jinqi) Chen <65606304+Kathryn-cat@users.noreply.github.com>
Some python tests are failing in the wheel. This PR skips them if the environment variable `WHEEL_TEST` is set.
This PR is related to https://github.com/tlc-pack/tlcpack/pull/115.
* Implement avg_pool2d slice op
* Address review comments and fix the STIR schedule
* Fix formatting issues
* Address pylint errors
* Additional formatting issues
* more pylint fixes
* Changed arch version to v68 for now
* Changing arch version back to v69
* Move the test to tests/python/contrib/test_hexagon/topi
The tvmc command will throw a error while the passed path of model is not exist, But for PaddlePaddle model, it contains 2 file model_name.pdmodel and model_name.pdiparams, we only pass the prefix like inference_model/model_name.
This pr is same with https://github.com/apache/tvm/pull/11108
Since the origin PR didn't update for a long time, I send this new PR
This PR enables task filtering in Apply-History-Best, which is used in
Relay/Relax integration. Previously, even though a task is ruled out
during task extraction, it still shows up in Relay compilation due to
the lack of filtering on `Apply-History-Best`. However, TE-to-TIR
conversion `te.CreatePrimFunc` doesn't support all cases with hybrid
operators involved, which leads to post-tuning failure affecting
multiple models.
* - add support for host-driven AoT execution on zephyr;
- add initial version of reference counting to prevent python code from inadvertently freeing tensors during garbage collection;
- add support for numerical indices to host-drive AoT get_input();
- add two initial tests for host-driven AoT execution on zephyr;
- rename existing zephyr AoT exec. test;
* address PR feedback
* increase stack size to accommodate qemu_riscv64 stack usage
This PR adds `te.extern_primfunc` which provides the interface around TE ExternOp that allows a TVMScript defined schedulable TIR PrimFunc to be inlined into a TE compute graph. The result is that TIR can be used for compute definitions in Relay OpStrategies and, paired with meta-scheduler support in relay as introduced in #10578, these compute definitions can be scheduled and tuned as demonstrated in the attached tests.
Prior to this, compute definitions were limited to those definable in TE only. As a consequence of this patch and ongoing improvements to TVMScript meta-programming (#11097), TOPI can be extended to include compute and scheduling functions targeting schedulable TIR uniformly.
* [BYOC] Make CUTLASS BYOC integration 'Collage friendly'
(See https://discuss.tvm.apache.org/t/byoc-supporting-cutlass-byoc-with-collage/12796/6 for
context, which in turn is part of Collage (https://github.com/apache/tvm-rfcs/blob/main/rfcs/0062-collage.md).
Currently CUTLASS has four entry points:
- The usual 'partition_for_cutlass' partitioning function, using the
standard pattern table and pass machinery (see cutlass/build.py).
- A 'tune_cutlass_kernels' function which augments CUTLASS partition
functions with the results of building and running test kernels (see cutlass/build.py).
- A 'relay.ext.cutlass' external codegen function which inspects the
turning results and generates a CSourceModule for each partitions
(see cutlass/codegen.cc).
- A 'build_cutlass_kernels_vm' function which runs 'export_library' with
all the nvcc compiler options needed to build all the CSourceModules
(see cutlass/bild.py).
For Collage we'd like CUTLASS to have only two entry points: 'partition_for_cutlass',
and 'relay.ext.cutlass' or equivalent. This makes the CUTLASS external codegen integration
composable with other integrations, which in turn helps Collage avoid having to understand any
external codegen APIs other than the global pattern table and the custom compilation function/pass.
Collage also tends to end up requiring multiple partitions for the same backend since it is
more aggressive at mixing-and-matching smaller sub-graphs between backends. Thus we'd also like
to make sure all tuning, generated code and compilation overhead is shared between all such CUTLASS
partitions.
So, in this PR:
- We add all the CUTLASS-specific tuning and compilation options as new Target
attributes for the 'external codegen' "cutlass" TargetKind (cutlass/target.cc).
The user now has one place to provide those settings, and we've already done the
legwork to plumb the target instance.
- We replace 'relay.ext.cutlass' with a 'RelayToTIR' custom pass hook
'CompileForCutlass' (see cutlass/codegen.cc). This pass obviously can see all
the CUTLASS partitions in the IRModule, so we can now share tuning results
between them all and can be sure to generate a single CSourceModule. The pass can
also invoke the compiler to yield a StaticModule, which we've also already done the
legwork to support. In this way all CUTLASS-specific steps are handled at once.
- For convenience we supply 'finalize_modules' and 'finalize_modules_vm' which
invoke nvcc for final linking (using export_library as usual). However, there's now
nothing CUTLASS specific in those helpers other than their overriding of the 'compiler' to
be nvcc.
- test_cutlass.py is updated to use the new API.
Though this is a breaking change for existing users of the CUTLASS integration the
change is pretty minor, as shown in test_cutlass.py.
* - Masa's comments
* - Remove unnecessary save.
In order to build a dataset for improving the cost model for MetaSchedule, I added several files
including importing models to TVM, extracting tuning tasks, and sampling measure candidates.
Meanwhile, I exposed some methods in C++ to the Python side to assist the process.
* unifies all MKLDNN/DNNL_CODEGEN to DNNL
* translate -lib=mkldnn to -libs=dnnl in target
* type check added before
* rebase and update conv2d from mkldnn to dnnl
* Reuse hexagon launcher in test session
* separate random name generation
* revert get_aot_executor
* Fix launcher for simulator case
* add stop server for simulator
* enable oneDNN conv op by using -libs=mkldnn
* add channel last format support and let oneDNN chose blocked format.
* remove unnecessary changes
* reformat 3 files
* reformat 1 file
* change the argument name
* change the argument name
* rename the arguments
* fix cpp lint issue
* fix cpp lint issue
* fix cpp lint issue
* clang reformated
* adjust .py import for testing
* function existence check in test
Per discussion with @Kathryn-cat
- [x] Move `initialize_with_tune_context` as private API `_initialize_with_tune_context`, and
encourage using `TuneContext.initialize`
- [x] Instead of using bunch of import statements, encourage using `ms.xxx` as the prefix
(e.g. `ms.database.MemoryDatabase`) to organize things better
- [x] Move `DefaultLLVM`, `DefaultCUDA` to a separate file and make them more discoverable
- [x] Move `DummyDatabase` to `tvm.meta_schedule.database.MemoryDatabase` given it's actually useful
- [x] Delegate class members' methods in `TuneContext`, for example, having
`TuneContext.generste_design_space` from `TuneContext.space_generator.generste_design_space`
Next PR:
- Allow using a string `"default"` in `TuneContext` as well as `tune_relay/tir/te` to quickly
specify a set of target-specific rules
- Add `TuneContext.tune` to allow directly tuning without task scheduler.
- Enhance detection of `ScheduleFn` in `TuneContext` to make it easier for users to quickly try out
template-driven scheduling on TIR.
Co-Authored-By: Kathryn (Jinqi) Chen <65606304+Kathryn-cat@users.noreply.github.com>
* [TIR, CUDA] Add pass to replace global to shared memory copy with cp.async
* add missing doc
* black
* missing src
* clang format
* clang format
* check against nested async scope
* [TVMSCRIPT] Improve tvmscript type hints
- Change numeric types to classes so they work as function arguments.
- Add var as a class.
- Add floordiv, index, and mod to PrimExpr.
* use Union
* [Schedule] Allowed string argument as block arg
This has previously been implemented for `Schedule.transform_layout`
in https://github.com/apache/tvm/pull/11296, extending to allow for
block arguments in all `Schedule` methods.
This change was only made for arguments that must be a `BlockRV`. For
arguments that may be either a `BlockRV` or another
type (e.g. `Schedule.get_child_blocks` accepts either `BlockRV` or
`LoopRV`), this sugar is not implemented, to avoid ambiguity.
* [Schedule] Allowed string argument to Schedule.reindex
Similar to https://github.com/apache/tvm/pull/11269, which added this
functionality to `Schedule.transform_layout`.
* CI test update
This PR fixes a few inconsistent pass registration and add testcases for them.
- `LowerInitBlock` had mismatch between its pass name and ffi key.
- `BindTarget`, `AnnotateEntryFunc`, `Filter` were not following the name convention of tir passes and they were not registered in FFI registry.
The CSE pass had been designed for potentially allowing comparisons (and commonings) of equivalent terms (like (x+y)+z and x+(y+z)), where **the notion of being equivalent was customizable, and no assumption was made about it**. That means that the implementation of the equivalence test function `EquivalentTerms()` - which was at the moment just calling the syntactical equality test `EqualTerms()` - could be replaced later by a cleverer equality test.
However, having such a generic way of comparing elements meant that in the function `SyntacticToSemanticComputations()`, where we were going from a hashtable of syntactical entities to what I called a vector of "semantical entites" (which are just canonical forms/representants of classes of equivalence of terms), **the only way was to compare each pair**.
That resulted in a quadratic behavior of this function, but there was no way around it as in order to merge equivalent entities into their class of equivalence, we had to compare them.
**This PR essentially does the following:**
- When computing the classes of equivalences of terms (therefore transforming a ComputationTable (i.e. a hashtable) into a vector of classes of equivalence) : **instead of comparing each pair of terms, relies on a normalization procedure to obtain a normal form for each of them**.
That transforms a small part of the algorithm that was quadratic to n.logn. However, it's difficult to see improvements in practice, in particular for average sized programs, as that part was a "small" quadratic to a "big" n.logn (finding things in a hash-table, copying it to a vector, etc).
It was probably going from a complexity of ~O(((n²-n)/2) + n.logn) to a complexity of ~O(3n + n.logn), so potential gains would only be expected for very large programs.
- Completely gives the user the possibility to turn ON/OFF the semantical comparisons of terms. It is turned OFF by default (as it's quite longer to compile with it ON, unsurprisingly), which means that by default, the equivalence coincides with the (syntactical) equality of terms.
As the pass was written with the possibility to do these additional commonings (like (x+y)+z and x+(y+z)), it was a good time to fully plug that completely, up to the Python user who can now turn that ON if he wants to. But again, it is OFF by default, so no real change on that.
To run it ON, simply do:
`with tvm.transform.PassContext(config={'tir.enable_equiv_terms_in_cse_tir':True}):`
before calling `build()`
- When this boolean is set to ON, it uses a simple implementation of the normalization function with equivalences that uses `arith::Analyzer::Simplify` as noted by in https://github.com/apache/tvm/pull/10544 . Note that this is not a real normalization procedure as it is incomplete (i.e., it is not guarantee to converge to the normal form), but it is correct, and it works well with most properties : associativity of +, distributivity of * on +, etc.
- Clarifies and enhance the test base for the pass. In particular, it adds the tests that were written in https://github.com/apache/tvm/pull/10544 but which did not make it through.
- Also add the test ( https://github.com/AndrewZhaoLuo/TVM-Sandbox/blob/19284ddbd6bb28af61c0c2aa8bb334c5c53731a7/tir/test_inconsistent_tir_lowering.py#L1 ) demonstrating the (older) non-deterministic lowering and put it into a proper test, as I found it useful for making sure that this does not happen again. It has been copied from https://github.com/apache/tvm/pull/10663 and only slightly adapted (in particular for doing the comparison of hashes automatically instead of printing them and relying on a human to compare them).
* Enable layer normalization in DNNL byoc.
* Added unittest for layer norm and make code compatible after introducing TensorRequisite(PR-11345)
* Fix lint issue
* Fix clang format issue
If a non-default location for the build directory is used, e.g. set via TVM_LIBRARY_PATH
we need to provide the user a way to overwrite CONFIGS_JSON_DIR as well.
Mehrdad Hessar