* [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
* [PROFILER] Fix percent compute bound calculation
Somehow the runtime was dropped from the percent compute bound
calculation. Tolerances on the test we bumped a little bit higher to try
and catch mistakes like this in the future.
* forgot print
* banker round op added based off tutorial
* black'd onnx.py file
* retriggering CI with empty commit due to autoscheduler test failure
* removed youtube link in comments
* retriggering CI due to test failure that passed locally
* [TE] Optimized version of concatenation layer
1. Concat implemented using extern_op
2. New tests added.
3. Workaround to allow inline extern_op-s with other layers.
* *test fix
* test_any.py fix.
* test_forward.py from tensorflow fix.
* lint fix.
* Fixes after code review.
* New comment added.
* Lint fix.
* Another lint fix.
* Comments added.
* rebase issue fix.
* Restored previous state.
* Update after code review.
* After code review changes.
* lint review.
* Change strategy for cuda to fix tests.
* Rebase to main
* Comments changes after review.
* Some more comments fixes.
* One more error fix in comments.
* restart build
This PR allows `relay.backend.MetaScheduleExtractTask` to take an extra argument `filter_func` which filters out tasks that don't need tuning. The counterpart of AutoScheduler is `traverse_to_get_io_tensors`.
* Test with CMSIS build added
disabled conv2d_nhwc_dsp.arm_cpu for non integers workloads
added debugging feature to TempDirectory
* revert arm_cpu strategy changes
* Address Andrew comments
* change copy to include
* add cmsis_path only as project option
* simplify (x * 96) % 64 to (x * 32) % 64
* adapt merge mulmod opt for OffsetOf computation
* merge DetectIterMap and DetectIterMapPadded
* adjust related interfaces for IterMapLevel
* - check incompatible left paddings
- determine case like x % 16, x in [0, 5) to be non-surjective, since usages may treat the region extent as 16 by mistake.
- skip second round of rewrite when there is no padding
- fix some typo in comments
* rebase upstream
billishyahao