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.
In some situations, such as before measuring the candidates, the arguments `run_secs`, `target`, and `args_info` in `TuningRecord` are not required. Per this request, the new `TuningRecord` API now accepts arguments in the order of `trace, workload, run_secs, target, args_info` with the last three being optional. Note that some tests might fail due to the change of argument order, so they might need to be adjusted accordingly.
Per discussion with @Kathryn-cat, we realized that the current API
design could be verbose if we only want to tune a single task, in which
case a dummy task scheduler still needs to be established to supply
`EvolutionarySearch` with proper `CostModel` and `Database`. This PR
fixes this UX issue.
Configuration is a place to store extra information related to the
specific profiler run. Right now it is just the executor used and the
number of threads. The roofline analysis also adds peak flops and peak
bandwidth.
Optimizes a case where padding appears as a separate nn.pad operation followed by a qnn.conv2d. If possible, the nn.pad will be partitioned and offloaded together with the qnn.conv2d operation, as opposed to separately. As a fallback, both operations will be considered separately.
cc Mousius NicolaLancellotti ekalda manupa-arm
* [microNPU] Fix output mismatch in Leaky ReLU
All codegen tests have been running with a representative dataset
between 0,1 which masked an output mismatch in Leaky ReLU when compared
to TFLite kernels. This issue can be replicated by replacing the
representative dataset range with something like -1,1.
To fix this mismatch, we use the same implementation for calculating
LUT values as Vela which uses arithmetic constrained to quantized
values, rather than the previously used floating point calculations.
Change-Id: I0ed52215acd27722873be609271971b6fc4aaef1
* fix lint
Change-Id: Ica7de0c000ee015e79fe10985b2ec7a9b341861f
* fix lint again
Change-Id: I005d90ad248bfff7090f99d161eefbdc962cba48
I exposed the initialize() method for TuneContextNode on the C++ side and added a corresponding method to TuneContext class on the Python side, so that we do not need to call initialize_with_tune_context for every scheduling rule.
In TE, a unit loop could be introduced by fusing an empty list of loops on a stage. This PR adds its counterpart in TIR, while being a bit more explicit with a new schedule primitive which adds a unit loop without impacting any existing functionalities.
* [CI] Improved skip messages when using @tvm.testing.requires_*
Previously, the same message was given regardless of why a test
couldn't be run. This has been split up into separate checks for TVM
cmake options in `config.cmake`, enabled targets in `TVM_TEST_TARGETS`
environment variable, and checks for available hardware.
* Refactor to specify repeated feature marks, compile-only markers
* Fixed lint errors
* Import from contrib, not from a different import
* Removed use of requires_llvm() as a list of marks
* Corrected mark from requires_gpu to requires_cuda
* Adding missing "not"
* Added USE_CMSISNN as a requirement for corstone300.
* [BYOC] Two helper passes for external codegen using RelayToTIR custom pass machinery
(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).
For reasons explained in the above thread I'm moving CUTLASS to be IRModule-at-a-time external codegen
using a custom RelayToTIR pass instead of the traditional function-at-a-time external codegen using
a relay.ext.cutlass registered function. This means some of the rewriing done on-the-fly by LowerTEPass now
needs to be done by the custom pass directly. This PR supplies two passes which ease that burden:
- Before starting the CUTLASS-specific processing, make sure all "Compiler" attributed functions have
unique global definitions (ie are outlined). Though functions start in this form after BYOC partitioning,
under Graph and AOT compilation flows those functions are then inlined to pass through the 'codegen' keyhole
which assumes the whole model is just one self-contained main function. This pass will undo that. (I gave up
trying to just remove the inlining in the first place.)
- After the CUTLASS-specific processing the now compiled "Compiler" attributed functions need to marked as
'extern'. The te_compiler.cc uses the "ExternalSymbol" attribute for that, but since a) the symbol name
is never needed, on the presense of the attribute is significant downstream and b) "ExternalSymbol" is
easy to confuse with "global_symbol", I just replaced "ExternalSymbol" with "Extern" with an Integer(1)
(cf "Primitive").
The outlining pass is a little more general than necessary because it (will also) be used by Collage to
rewrite the IRModule into optimally partitioned form while making maximal reuse of partition functions.
Hence the abstract GlobalSymbolCache.
* - Andrew's comments
Jason