[Vulkan] Add im2col wmma implementation for CHW direct conv2d - #15002
Closed
mei-ye wants to merge 51 commits into
Closed
[Vulkan] Add im2col wmma implementation for CHW direct conv2d#15002mei-ye wants to merge 51 commits into
mei-ye wants to merge 51 commits into
Conversation
Add SPIR-V code generation for "SPV_NV_cooperative_matrix" extension. Add a matrix multiplicaiton unit test.
apache#14856) [TOPI] Fix in interpretation of empty axis parameter in reduction functions
* fix the undeclared identifier 'f' * Update codegen_cpu.cc
This PR enahnces create_staticlib to take tar files which can be handy when combining multiple libs into a single a file.
* [AutoTVM] Use f-strings for string formatting, NFC Replace uses of % and .format() with f-strings. Reformat modified files. * Fix typo: epochs -> epoch
This PR enhances systemlib to make it unique per prefix. Can help reduce flaky problems in multiple loading of same lib.
A separate channel-dimension nn.pad relay operator is rewritten as Relay concatenate operation. --------- Co-authored-by: Sergey Smirnov <89378719+sergey-grovety@users.noreply.github.com> Co-authored-by: arina.naumova <naumova@grovety.com>
* fix broadcast_tensors * Update pytorch.py * Update test_forward.py * Update test_forward.py
* [CI] Update CPU image to install PyTorch * use link-static to prevent symbol conflict problem * tlcpackstaging -> tlcpack * disable tvmc pth tests for now * fixed skip marker
* [TVMScript] Cherry-pick module.other_func syntax from unity * [TIR][TVMScript] Added format/parsing of subroutine calls Similar to `module.relax_func(args)` syntax used when parsing Relax functions, this allows `module.tir_func(args)` to be used when parsing TIR PrimFuncs.
* [Contrib] Use f-strings for string formatting, NFC Replace uses of % and .format() with f-strings. Reformat modified files. * Fix linter
* [TIR] Expand unit tests for ConvertSSA Prior to this PR, there was a single test which invoked ConvertSSA and checked that no error was thrown. This PR adds additional test cases for nested variable definition, and for variables de-duplicated across separate function calls in an `IRModule`. Of the additional tests, the behavior tested by `TestDedupAutoBroadcastBuffer` and `TestReusedBufferParameter` fails on main, and is resolved by this PR. * Update tests to avoid relying on TVMScript output of non-SSA The "before" cases must have SSA violations, which is not valid in TIR. As such, future versions of the TVMScript parser may remove the SSA violations in the process of parsing. Updated unit tests introduce SSA violations through the Python API, to avoid this potential breakage.
…#14900) Previously, a `DeclBuffer` object in which `decl_buffer->buffer->data` had not been defined would be printed without the `data` argument. This was the same representation as the `Allocate`/`DeclBuffer` pattern (e.g. `buf = T.decl_buffer(shape,dtype)`), and so the `Allocate` node would be inserted when parsing. This commit updates the printing of buffers to be aware of which variables are implicitly declared in the process of declaring a buffer. (e.g. A buffer that occurs in `BlockNode::alloc_buffers` defines both the buffer and the data pointer, while a `DeclBuffer` only defines the buffer.) This is used to produce a different representation for the undefined data pointer, removing the ambiguity with the `Allocate`/`DeclBuffer` sugar. Because this change only affects malformed TIR in which a `DeclBuffer` node references an undefined data pointer, it is intended primarily for debugging purposes.
…#14901) * [CodeGen][LLVM] Codegen to generate internal functions Previously, `CodeGenLLVM` required all TIR PrimFuncs to have the `kGlobalSymbol` attribute, using its value as the externally-visible symbol in the generated library. This commit relaxes that requirement, using the presence of `kGlobalSymbol` to indicate whether a function should be exposed externally. If `kGlobalSymbol` is not defined, then the symbol name is generated from the name of the `tvm::GlobalVar` with the prefix `"_internal_"`, and the symbol is not exposed externally. Since this does not change the codegen behavior for any function that was previously supported, this is not a breaking change. * [Codegen][LLVM] Handle callsite for internal functions * [UnitTest][LLVM] Added test for LLVM codegen for subroutine
… all PyTorch versions (apache#14903) * fix test cases Add some version check to avoid calling the operator/attribute which has not been introduced yet！ * Update test_forward.py
* fix a wrong crash message * Update keras.py
* add a dimension check * Update test_forward.py * Update pytorch.py * Update pytorch.py
[Docker] Add polly package. llvm-16 packages did not include the polly library, but llvm-config expects it to be there, so tvm failed to build. This change adds libpolly-16-dev to the ci_cpu package
…pache#14934) fix unreachable test cases
…or with a stride > 3 (apache#14861) The nn.avg_pool2d operator with a stride size greater than 3 in any of the spatial dimensions is rewritten as ethosu avg_pool with strides=[1,1] if the case satisfies the additional conditions (no AvgPool2D padding, spatial dimensions of ifm and shape of pooling are equal). --------- Co-authored-by: Sergey Smirnov <89378719+sergey-grovety@users.noreply.github.com> Co-authored-by: arina-grovety <> Co-authored-by: Arina Naumova (grovety.com) <naumova@grovety.com> Co-authored-by: Arina <117634809+arina-grovety@users.noreply.github.com>
…ache#14915) [TVMScript] Allow T.target("device", host="host") in TVMScript Prior to this commit, the `TargetNode::host` could be specified in TVMScript as part of the config dictionary, under the key `"host"`. However, this required all other device parameters to be explicitly specified, rather than using any of the short-hand string representations. This commit forwards the `host` argument from TVMScript's `T.target` method to `tvm.target.Target`, allowing both the device and host to be specified using the shorthand string representation. ```python @T.prim_func def before_this_commit(): T.func_attr( { "target": T.target( { "arch": "sm_86", "host": {"keys": ["cpu"], "kind": "llvm", "tag": ""}, "keys": ["cuda", "gpu"], "kind": "cuda", "max_num_threads": 1024, "tag": "", "thread_warp_size": 32, } ) } ) T.evaluate(0) @T.prim_func def after_this_commit(): T.func_attr({"target": T.target("cuda", host="llvm")}) T.evaluate(0) ```
This PR enables log before throw in windows. The error handling in windows sometimes can be tricky and it is helpful to be able to get as many information as possible. Turning this option on can help us detect possible errors here.
* [Target] Add target to all TVM callbacks This PR adds an extra parameter `target` to all `tvm_callback_*` so that the callback can decide its own behavior by querying which target to compile against. * fix lint * fix lint
A small script that exposes host CPU name, target triple and features:
<details>
```python
import tvm
def main():
get_default_target_triple = tvm._ffi.get_global_func("tvm.codegen.llvm.GetDefaultTargetTriple")
get_process_triple = tvm._ffi.get_global_func("tvm.codegen.llvm.GetProcessTriple")
get_host_cpu_name = tvm._ffi.get_global_func("tvm.codegen.llvm.GetHostCPUName")
get_host_cpu_features = tvm._ffi.get_global_func("tvm.codegen.llvm.GetHostCPUFeatures")
target_triple = get_default_target_triple()
process_triple = get_process_triple()
host_cpu_name = get_host_cpu_name()
host_cpu_features = get_host_cpu_features()
print("target_triple: {}".format(target_triple))
print("process_triple: {}".format(process_triple))
print("host_cpu_name: {}".format(host_cpu_name))
print("host_cpu_features:")
for name, value in host_cpu_features.items():
print(" {}: {}".format(name, bool(value)))
if __name__ == "__main__":
main()
```
</details>
Output (AMD CPU):
<details>
```
target_triple: x86_64-unknown-linux-gnu
process_triple: x86_64-unknown-linux-gnu
host_cpu_name: znver2
host_cpu_features:
xsaveopt: True
tsxldtrk: False
sse: True
movdiri: False
mmx: True
pku: False
amx-int8: False
amx-tile: False
rdpid: True
avx512vbmi2: False
cmov: True
widekl: False
f16c: True
bmi: True
gfni: False
avx512cd: False
movdir64b: False
rdseed: True
clwb: True
avx512er: False
avx512f: False
sse4.2: True
avxifma: False
sse2: True
avx512vp2intersect: False
prfchw: True
avx512pf: False
vaes: False
waitpkg: False
amx-bf16: False
prefetchi: False
uintr: False
fxsr: True
bmi2: True
lzcnt: True
avx512vbmi: False
avx512bf16: False
prefetchwt1: False
xsaves: True
movbe: True
rtm: False
pclmul: True
hreset: False
sahf: True
fma4: False
xop: False
vpclmulqdq: False
sgx: False
avx512vnni: False
popcnt: True
xsavec: True
aes: True
avx512vpopcntdq: False
kl: False
avx512bitalg: False
xsave: True
avxvnni: False
raoint: False
clflushopt: True
sse4a: True
avx512bw: False
cx16: True
avxvnniint8: False
amx-fp16: False
cldemote: False
rdrnd: True
ptwrite: False
rdpru: True
avx: True
adx: True
avx512vl: False
pconfig: False
shstk: False
64bit: True
crc32: True
sha: True
cmpccxadd: False
tbm: False
serialize: False
mwaitx: True
avx512ifma: False
avx512fp16: False
clzero: True
avx2: True
cx8: True
fma: True
lwp: False
enqcmd: False
wbnoinvd: True
sse4.1: True
avx512dq: False
ssse3: True
fsgsbase: True
invpcid: False
sse3: True
avxneconvert: False
```
</details>
Note that LLVM doesn't guarantee automatic feature detection always succeeds, particularly for newer CPU models and older LLVM builds (e.g. M2 CPU + LLVM 16), the result is usually inaccurate. In this case, i.e. CPU feature detection fails, we will print a warning message and return an empty dict instead.
To properly detect CPU features on macbook, the commands below provided by the system are the most accurate:
```bash
sysctl -a machdep.cpu
sysctl -a hw.optional
```
On linux, usually it is recommended to directly query via:
```bash
cat /proc/cpuinfo
```
* fix conv transpose import from TF * fix String::fromwe() to String::from() * torch squeeze can use a list of axis * added test for squeeze with multiple axis (pytorch 2) * clean old code * code without reformating --------- Co-authored-by: Mikael Sevenier <mikael.sevenier@sima.ai>
* [TIR] MakePackedAPI, handle missing kGlobalSymbol Previously, `MakePackedAPI` required all functions to have the `kGlobalSymbol` attribute. This commit updates the behavior such that `MakePackedAPI` only modifies PrimFuncs that have the `kGlobalSymbol` attribute, and passes through any other PrimFunc unmodified. * [TIR] Update calls to externally-exposed subroutines in MakePackedAPI When a function is updated to use the `PackedFunc` API, any calls made to that function from elsewhere in the `IRModule` should be updated as well. * Bugfix, don't update the callsite unless the callee is also updated
Prior to this commit, MakeUnpackedAPI required all functions to be annotated with `kGlobalSymbol` (`"global_symbol"`). This commit updates the transformation to apply only to functions that have the `kGlobalSymbol` attribute, and to update any internal callsites of the modified functions. This is analogous to the changes made in apache#14913, which updates `MakePackedAPI`.
…ache#14945) * [Target] Added utility method TargetNode::HasKey() This utility method makes it easier to determine if a target contains a specific key. * [TIR] Added utility method tvm::tir::IsHostFunc(const PrimFunc&) For modules that contain both host and device functions, this utility function checks whether a given PrimFunc is a host function, based on the target annotation. * [TIR] Restrict tir.transform.CombineContextCall to host functions Previously, the `tir.transform.CombineContextCall` pass applied to all functions in an `IRModule`, but was only applied to modules that contain only host functions. This commit updates `tir.transform.CombineContextCall` to apply only to host functions.
…apache#14950) PRs apache#14913 and apache#14914 made analogous changes to `MakePackedAPI` and `MakeUnpackedAPI` to handle subroutine calls. Both PRs introduced the same symbol, `tvm::tir::SubroutineCallRewriter`, a local utility to update internal calls to a modified function. While each PR passed CI individually, and was therefore able to merge, having both changes caused a duplicate symbol. This commit updates `MakePackedAPI` and `MakeUnpackedAPI` to place their local utilities into anonymous namespaces, avoiding the conflict.
…che#14942) * [Bugfix] Avoid symbol conflicts in MakePackedAPI/MakeUnpackedAPI PRs apache#14913 and apache#14914 made analogous changes to `MakePackedAPI` and `MakeUnpackedAPI` to handle subroutine calls. Both PRs introduced the same symbol, `tvm::tir::SubroutineCallRewriter`, a local utility to update internal calls to a modified function. While each PR passed CI individually, and was therefore able to merge, having both changes caused a duplicate symbol. This commit updates `MakePackedAPI` and `MakeUnpackedAPI` to place their local utilities into anonymous namespaces, avoiding the conflict. * [Target] Added WithoutHost method * [TIR] Preserve existing kTarget function attribute in BindTarget Previously, if a function already has a `tvm::attr::kTarget` attribute, it will be overwritten by the `tir.BindTarget` transform. This commit updates the behavior such that `tir.BindTarget` adds annotations to functions that are missing a target annotation, but preserves any existing target annotations. This is part of a series of commits to simplify the handling of multi-target builds.
…che#14943) * [Bugfix] Avoid symbol conflicts in MakePackedAPI/MakeUnpackedAPI PRs apache#14913 and apache#14914 made analogous changes to `MakePackedAPI` and `MakeUnpackedAPI` to handle subroutine calls. Both PRs introduced the same symbol, `tvm::tir::SubroutineCallRewriter`, a local utility to update internal calls to a modified function. While each PR passed CI individually, and was therefore able to merge, having both changes caused a duplicate symbol. This commit updates `MakePackedAPI` and `MakeUnpackedAPI` to place their local utilities into anonymous namespaces, avoiding the conflict. * [Target] Added utility method TargetNode::HasKey() This utility method makes it easier to determine if a target contains a specific key. * [TIR] Added utility method tvm::tir::IsHostFunc(const PrimFunc&) For modules that contain both host and device functions, this utility function checks whether a given PrimFunc is a host function, based on the target annotation. * [TIR] Restrict InstallDebugSpans to host functions Previously, the `tir.InstallDebugSpans` pass required the module to contain only a single PrimFunc. This commit relaxes the requirement, to require a single host-side PrimFunc, and to ignore any other device-side functions.
…che#14941) * [Bugfix] Avoid symbol conflicts in MakePackedAPI/MakeUnpackedAPI PRs apache#14913 and apache#14914 made analogous changes to `MakePackedAPI` and `MakeUnpackedAPI` to handle subroutine calls. Both PRs introduced the same symbol, `tvm::tir::SubroutineCallRewriter`, a local utility to update internal calls to a modified function. While each PR passed CI individually, and was therefore able to merge, having both changes caused a duplicate symbol. This commit updates `MakePackedAPI` and `MakeUnpackedAPI` to place their local utilities into anonymous namespaces, avoiding the conflict. * [TVMScript] Prevent bool to int conversion in T.Assert condition Previously, while literal `True` and `False` values were converted to `tvm::Bool` instances, constant-foldable expressions (e.g. `0 == 1`) would be evaluated to `True`, but were then passed directly to the FFI. Because the FFI uses the same representation for integer and boolean values, the conversion to `PrimExpr` resulted in a `tvm::Integer` instead of `tvm::Bool`. This commit converts the argument of `T.Assert` to a `tvm::Bool` before calling the FFI, avoiding the ambiguity. In addition, the `AssertStmt` constructor now validates the datatype of the condition, to prevent it from re-occurring. This was first caught in the unit test `test_debug_info.py::test_llvm_ir_debug_info`, which failed on some versions of LLVM due to the use of `i32` as the condition of an assert. * Updated TVMScript printer unit test with boolean condition
* [Bugfix] Avoid symbol conflicts in MakePackedAPI/MakeUnpackedAPI PRs apache#14913 and apache#14914 made analogous changes to `MakePackedAPI` and `MakeUnpackedAPI` to handle subroutine calls. Both PRs introduced the same symbol, `tvm::tir::SubroutineCallRewriter`, a local utility to update internal calls to a modified function. While each PR passed CI individually, and was therefore able to merge, having both changes caused a duplicate symbol. This commit updates `MakePackedAPI` and `MakeUnpackedAPI` to place their local utilities into anonymous namespaces, avoiding the conflict. * [Codegen][LLVM] Allow void return type from PackedFunc Previously, calling a packed func that returns void would result in an error being raised from `tir::APIType`, as there is no runtime representation of a void type. This commit updates `CodeGenCPU::MakeCallPackedLowered` to only read the return value and type fo a `PackedFunc` when the TIR return type is non-void.
Current codegen output `(half4)*(device uint*)A` tries to create a `int32` number and then cast it to `half4`, which is not the expected behavior. As Metal supports `uchar4` and `char4` types, we can direct use them to solve that problem.
This PR refactors SplitHostDevice into three separate transformations. Previously, SplitHostDevice would replace device regions with a builtin::tvm_call_packed() node to replace the extracted region. After this PR, this process is performed in three separate steps.
AnnotateDeviceRegion: Annotate the regions that should be executed on another target.
SplitHostDevice: Extract the annotated region into an independent PrimFunc, with a GlobalVar to represent the call from into the new subroutine.
LowerDeviceKernelLaunch: For any subroutine call where the caller and callee are on different devices, replace with a device kernel launch.
* PR#14915 [TVMScript] Allow T.target("device", host="host") in TVMScript
Prior to this commit, the `TargetNode::host` could be specified in
TVMScript as part of the config dictionary, under the key `"host"`.
However, this required all other device parameters to be explicitly
specified, rather than using any of the short-hand string
representations. This commit forwards the `host` argument from TVMScript's
`T.target` method to `tvm.target.Target`, allowing both the device and
host to be specified using the shorthand string representation.
```python
@T.prim_func
def before_this_commit():
T.func_attr(
{
"target": T.target(
{
"arch": "sm_86",
"host": {"keys": ["cpu"], "kind": "llvm", "tag": ""},
"keys": ["cuda", "gpu"],
"kind": "cuda",
"max_num_threads": 1024,
"tag": "",
"thread_warp_size": 32,
}
)
}
)
T.evaluate(0)
@T.prim_func
def after_this_commit():
T.func_attr({"target": T.target("cuda", host="llvm")})
T.evaluate(0)
```
* [Target] Added WithoutHost method
* [TIR] SplitHostDevice, handle missing kGlobalSymbol
Previously, the symbol name of the extracted compute kernel was
defined based on the `kGlobalSymbol` attribute, which was required to
be present. This commit updates `SplitHostDevice` to generate the
symbol name using `kGlobalSymbol` if present, and to fall back to the
name of the `tvm::GlobalVar` for internal functions.
* [TIR] Refactor SplitHostDevice into three separate passes
First pass, `AnnotateDeviceRegions`. This pass decides which portions
of a PrimFunc should be run on the device, and annotates them with
`kTarget` attribute, indicating which target should be used for later
lowering steps.
Second pass, `SplitHostDevice`. This pass extracts the annotated
region into an independent PrimFunc. The `kTarget` attribute of the
extracted kernel is defined by the `kTarget` annotation inserted by
`AnnotateDeviceRegions`. The host function is marked by the
`tvm::tir::attr::kIsHostFunc` attribute, allowing it to be recognized
by later host-only lowering passes.
Third pass, `LowerDeviceKernelLaunch`. This pass identifies
subroutine calls that call into device kernels, and rewrites them into
`T.tvm_call_packed`.
* Add unit tests specifically for SplitHostDevice behavior
* Added unit test specifically for AnnotateDeviceRegions
* Added unit tests for LowerDeviceKernelLaunch
* Minor cleanup, moved all kernel launch collection into one spot
Previously, the SplitHostDevice pass added the
`tir::attr::kKernelLaunchParams` attribute, and the
LowerDeviceKernelLaunch pass filled in the values for it. This
cleanup makes the kernel launch params be the sole responsibility of
LowerDeviceKernelLaunch.
* Updated unit tests for LowerWarpMemory
* Updated unit tests for ThreadSync
* Updated unit test for inject ptx async copy
* [Bugfix] Avoid symbol conflicts in MakePackedAPI/MakeUnpackedAPI
PRs apache#14913 and
apache#14914 made analogous changes to
`MakePackedAPI` and `MakeUnpackedAPI` to handle subroutine calls.
Both PRs introduced the same symbol,
`tvm::tir::SubroutineCallRewriter`, a local utility to update internal
calls to a modified function. While each PR passed CI individually,
and was therefore able to merge, having both changes caused a
duplicate symbol.
This commit updates `MakePackedAPI` and `MakeUnpackedAPI` to place
their local utilities into anonymous namespaces, avoiding the
conflict.
* Maintain "tir.is_global_func" attr in device-side entry point
* SplitHostDevice, update the host-side target to be the host
* [TIR] Update LowerDeviceKernelLaunch to avoid kIsHostFunc
Update to use the `tvm::tir::IsHostFunc` utility function, rather than
the `kIsHostFunc` attribute. Per discussion on
apache#14020, the `kIsHostFunct` attribute
should only be used in `BindTarget`, and should not be re-introduced
in `SplitHostDevice`.
* Remove is_host_func from SplitHostDevice tests
* [Runtime,RPC] Use f-strings for string formatting, NFC Replace uses of % and .format() with f-strings. Reformat modified files. * Fix linter
* [Hexagon] Add support for v73, make v68 default Snapdragon 8 Gen2 is v73, Hexagon SDK 5.2+ supports it. * Guard CPU versions with #ifdefs The CPU version constants are enums, but they are initialized with values that are #defined in HexagonVersions.h. Before using any HEX_CPU_Vxx, check if the corresponding initializer has been #defined.
* [TVMScript][TIR] Parse subroutine calls with no arguments In most cases, the IR dialect in `GlobalVar.__call__` can be inferred from the argument types. If there are no arguments, then the returned value is ambiguous. This commit updates the TIR parser to identify and fix this case of erroneously producing a `relay.Call` instead of `tir.Call`. In addition, to prevent this from re-occuring, an unrecognized type resulting from `def visit_expr_stmt` now results in an error, rather than being silently ignored. * Ignore str for unknown parser result These may are used as docstrings in the TVMScript, even though they are not represented in the TIR. * Lint fixes
Prior to this commit, `ArgBinder` would always introduce a new variable to represent the input argument, even if the argument already a primitive type. This introduces trivial let bindings that are expected to be simplified out, but which can produce dangling `tir::Var` usage in some cases (see apache#14951). This commit updates `ArgBinder` to prefer using the original `tir::Var` when possible. That is, when a function takes `n: T.int32` as input, the packed function should produce a binding `n: T.int32 = T.tvm_struct_get(...)`, rather than producing a binding `arg_n = T.tvm_struct_get(...)` followed by `n = arg_n`.
* [IR][SIBuilder] - Add SIBuilder to handle the span propagation between passes - Add SequentialSpan for multiple source expressions conversion between passes - Add test cases for SIBuilder and SequentialSpan * [IR][SIBuilder] - Make null implementation as base class - Add comments and change naming based on reviewing --------- Co-authored-by: Joey Tsai <chunit@qti.qualcomm.com>
…pache#14982) This resolves an issue introduced by the combination of apache#14918 and apache#14945. The bug occurred for targets that do not require device-side codegen, but do require a `device_type` other than `kDLCPU`. It wasn't caught by CI, as the issue only occurred with the combination of both PRs. 1. apache#14918 updated `SplitHostDevice` to only modify the `"target"` attribute when a device-side function has been extracted. 2. For VTA, there is no device-side function, as everything is done through host-side API calls. 3. From (1) and (2), the VTA examples kept the target `T.target("ext_dev", host="llvm")` after the `SplitHostDevice` pass, instead of being updated to `T.target("llvm")`. 4. apache#14945 restricted CombineContextCall to only apply to host-side passes. 5. From (4) and (5), the `CombineContextCall` pass was no longer applied to the VTA context calls. This PR fixes `SplitHostDevice`, updating the target from `T.target("ext_dev", host="llvm")` to `T.target("llvm")`, even if no device sections have been extracted from the function.
# The Issue
Currently, the ThreadAllReduce pass would throw an error when the mask data type is uint32 and the `group_index`'s data type is int64:
```bash
1: tvm::tir::ThreadAllreduceBuilder::MakeAllreduce(tvm::tir::CallNode const*)
at /home/zhye/repos/relax/src/tir/transforms/lower_thread_allreduce.cc:362
0: tvm::tir::BufferStore::BufferStore(tvm::tir::Buffer, tvm::PrimExpr, tvm::runtime::Array<tvm::PrimExpr, void>, tvm::Span)
at /home/zhye/repos/relax/src/tir/ir/stmt.cc:477
File "/home/zhye/repos/relax/src/tir/ir/stmt.cc", line 477
TypeError: dtype mismatch on BufferStore: buffer's dtype is `uint32`, the lanes of indexing are: `1`, but RHS's dtype is `int64`
```
As int64 becomes the standard index data type for large models, we should fix the issue.
# The Fix
This PR resolves the issue by casting the `group_index` to the data type used in mask.
This PR add a dummy target that can be renamed as libtvm and can be used to create a dependency on libtvm without compiling everything. This can be useful for down stream in platforms such as windows.
* [Hexagon] Update instructions to compile hexagon runtime Also, added include path of dlfcn.h which is applicable from Hexagon SDK 5.3.0.0 * Add condition to enable dlfcn path from SDK 5.3.0.0
Print error string if it is impossible to build kernel.
…14902) This PR updates the Module and Registry's DLL function to use String container instead of std::string. While it is impossible to obtain a stable ABI due to the nature of c++, and it is important to keep that flexibility, it is helpful to keep small set of tvm/runtime functions to work with use a String so it is more stable across compilers.
Currently only 16x16x16 FP16 wmma fragments with FP32 intermediates are supported. Add "min_design_space" as a parameter to give minimum design space for meta scheduler tuning. Add "use_int32_const" as a paramter to use int32 type for constants. Allow target query to be called from the schedules so that samplings are constrained to produce legal schedules. Add a unit test test_meta_schedule_vulkan_wmma.py
Collaborator
|
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.
Generated by tvm-bot |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
17 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Currently only 16x16x16 FP16 wmma fragments with FP32 intermediates are supported. Add "min_design_space" as a parameter to give minimum design space for meta scheduler tuning. Add "use_int32_const" as a paramter to use int32 type for constants. Allow target query to be called from the schedules so that samplings are constrained to produce legal schedules. Add a unit test test_meta_schedule_vulkan_wmma.py