[Hexagon] Remove "c" as aot_host_target tvm/contrib/hexagon/pytest_pl…

[kparzysz-quic]

…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.

[tvm-bot]

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.

• cc @ibsidorenko _{See #10317 for details}

_{Generated by tvm-bot}

[kparzysz-quic]

cc: @Lunderberg
Maybe you know why it worked before without spending time on investigating... The current situation isn't wrong, but I'm wondering what we did before that was different.

[Lunderberg]

Hmm. I think it might be related to this change 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.

[kparzysz-quic]

LowerIntrin inserted fma, which is a TIR builtin. Then, Hexagon code legalized it to an LLVM intrinsic.

Edit: Hexagon legalization assumes LLVM codegen, so it can unconditionally generate LLVM-specific code. This is what the PR comment refers to.

[Lunderberg]

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.