You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "kparzysz-quic (via GitHub)" <gi...@apache.org> on 2023/05/31 15:00:35 UTC
[GitHub] [tvm] kparzysz-quic opened a new pull request, #14997: [Hexagon] Remove "c" as aot_host_target tvm/contrib/hexagon/pytest_pl…
kparzysz-quic opened a new pull request, #14997:
URL: https://github.com/apache/tvm/pull/14997
…ugin.py
After PR#14918 (changes to SplitHostDevice), LowerIntrin code generates fma, which then becomes the LLVM intrinsic fmuladd. Problem is, the code then goes to the C codegen, which knows nothing about LLVM intrinsics. The result is an abort at compile-time:
```
E File ".../src/target/source/codegen_c.cc", line 611
E TVMError: Unresolved call Op(tir.call_llvm_pure_intrin)
```
Why this didn't happen before is unclear at the moment, but using C codegen with Hexagon code doesn't sound like a winning combination. Hexagon-related code assumes LLVM codegen, and the above crash is the result of that.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] kparzysz-quic commented on pull request #14997: [Hexagon] Remove "c" as aot_host_target tvm/contrib/hexagon/pytest_pl…
Posted by "kparzysz-quic (via GitHub)" <gi...@apache.org>.
kparzysz-quic commented on PR #14997:
URL: https://github.com/apache/tvm/pull/14997#issuecomment-1570511445
LowerIntrin inserted `fma`, which is a TIR builtin. Then, Hexagon code legalized it to an LLVM intrinsic.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tvm-bot commented on pull request #14997: [Hexagon] Remove "c" as aot_host_target tvm/contrib/hexagon/pytest_pl…
Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #14997:
URL: https://github.com/apache/tvm/pull/14997#issuecomment-1570405272
<!---bot-comment-->
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](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<!--bot-comment-ccs-start-->
* cc @ibsidorenko <sub>See [#10317](https://github.com/apache/tvm/issues/10317) for details</sub><!--bot-comment-ccs-end-->
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] quic-sanirudh merged pull request #14997: [Hexagon] Remove "c" as aot_host_target tvm/contrib/hexagon/pytest_pl…
Posted by "quic-sanirudh (via GitHub)" <gi...@apache.org>.
quic-sanirudh merged PR #14997:
URL: https://github.com/apache/tvm/pull/14997
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] kparzysz-quic commented on pull request #14997: [Hexagon] Remove "c" as aot_host_target tvm/contrib/hexagon/pytest_pl…
Posted by "kparzysz-quic (via GitHub)" <gi...@apache.org>.
kparzysz-quic commented on PR #14997:
URL: https://github.com/apache/tvm/pull/14997#issuecomment-1570409667
cc: @Lunderberg
Maybe you know why it worked before without spending time on investigating...
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Lunderberg commented on pull request #14997: [Hexagon] Remove "c" as aot_host_target tvm/contrib/hexagon/pytest_pl…
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on PR #14997:
URL: https://github.com/apache/tvm/pull/14997#issuecomment-1570486985
Hmm. I think it might be related to [this change](https://github.com/apache/tvm/pull/14918/files#diff-c795c8cdf29dcfda3dcb7bf342247b22deee71e2a9c037046254625e91bdd152L264) in particular. Previously, the host function had its `kTarget` attribute removed altogether. Afterwards, the host function had its `kTarget` attribute updated to be the original functions `target->GetHost()`. Effectively, that would result in any per-function host annotations being replaced with the host passed to `tvm.build`, which defaults to LLVM.
That still wouldn't explain how the LLVM intrinsics were inserted in the first place, as `LowerIntrin` should use whatever the function annotation is at the point of the pass being applied, but it might be close to an answer.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Lunderberg commented on pull request #14997: [Hexagon] Remove "c" as aot_host_target tvm/contrib/hexagon/pytest_pl…
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on PR #14997:
URL: https://github.com/apache/tvm/pull/14997#issuecomment-1570645910
Ah, got it. I thought that the LLVM fma was being directly inserted by `tir.transform.LowerIntrin`.
Summing up, it sounds like Hexagon legalization assumes LLVM codegen, and so the `"c"` codegen shouldn't have been used in this test case. The test case worked previously, because `SplitHostDevice` removed the host information entirely, then defaulted to LLVM. The refactor in #14918 fixed that behavior in `SplitHostDevice`, which meant that it was no longer covering up the use of `"c"` runtime.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org