This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 8f40544c85 [LLVM] Fix insertDbgValueIntrinsic for Metal backend
(#18706)
8f40544c85 is described below
commit 8f40544c85c7f3b974091728748c6885b05895e4
Author: Ruihang Lai <[email protected]>
AuthorDate: Mon Feb 2 08:52:10 2026 -0800
[LLVM] Fix insertDbgValueIntrinsic for Metal backend (#18706)
Some TIR code could not be compiled to Metal due to an LLVM issue, as
described in #18585. This PR fixes the issue according to the Option 2
in
https://github.com/apache/tvm/issues/18585#issuecomment-3649857591.
Unit tests are updated, though they are not enabled in CI for now.
---
src/target/llvm/codegen_llvm.cc | 26 ++++++++++++++++++++++
tests/python/codegen/test_gpu_codegen_allreduce.py | 5 ++++-
tests/python/codegen/test_target_codegen_metal.py | 2 +-
3 files changed, 31 insertions(+), 2 deletions(-)
diff --git a/src/target/llvm/codegen_llvm.cc b/src/target/llvm/codegen_llvm.cc
index 131c8212c5..ed2f927c0d 100644
--- a/src/target/llvm/codegen_llvm.cc
+++ b/src/target/llvm/codegen_llvm.cc
@@ -2323,6 +2323,32 @@ void CodeGenLLVM::AddDebugInformation(llvm::Value*
llvm_value, const Var& tir_va
auto* di_loc = llvm::DILocation::get(*llvm_target_->GetContext(), 0, 0,
di_subprogram_);
+#if TVM_LLVM_VERSION >= 150
+ // LLVM 15+ requires dbg_declare to reference pointer or integer types only.
+ // For non-pointer types (floats, vectors), use dbg_value instead to track
+ // the SSA value directly rather than a memory location.
+ if (!llvm_value->getType()->isPointerTy()) {
+ if (insert_before) {
+ // LLVM 20+ changed insertDbgValueIntrinsic to take BasicBlock::iterator
+ // instead of Instruction* for the insertion point.
+#if TVM_LLVM_VERSION >= 200
+ dbg_info_->di_builder_->insertDbgValueIntrinsic(
+ llvm_value, local_var, dbg_info_->di_builder_->createExpression(),
llvm::DebugLoc(di_loc),
+ llvm::BasicBlock::iterator(insert_before));
+#else
+ dbg_info_->di_builder_->insertDbgValueIntrinsic(llvm_value, local_var,
+
dbg_info_->di_builder_->createExpression(),
+ llvm::DebugLoc(di_loc),
insert_before);
+#endif
+ } else {
+ dbg_info_->di_builder_->insertDbgValueIntrinsic(
+ llvm_value, local_var, dbg_info_->di_builder_->createExpression(),
llvm::DebugLoc(di_loc),
+ builder_->GetInsertBlock());
+ }
+ return;
+ }
+#endif
+
if (insert_before) {
#if TVM_LLVM_VERSION >= 200
dbg_info_->di_builder_->insertDeclare(
diff --git a/tests/python/codegen/test_gpu_codegen_allreduce.py
b/tests/python/codegen/test_gpu_codegen_allreduce.py
index 31b6511e0e..5d73fb36bb 100644
--- a/tests/python/codegen/test_gpu_codegen_allreduce.py
+++ b/tests/python/codegen/test_gpu_codegen_allreduce.py
@@ -15,6 +15,7 @@
# specific language governing permissions and limitations
# under the License.
import tvm
+import tvm_ffi
import tvm.testing
import numpy as np
from tvm.script import tir as T
@@ -96,7 +97,9 @@ def
optional_metal_compile_callback(define_metal_compile_callback):
@tvm.register_global_func(name, override=True)
def compile_metal(src, target):
- return tvm.contrib.xcode.compile_metal(src, sdk="macosx")
+ from tvm.contrib.xcode import compile_metal # pylint:
disable=import-outside-toplevel
+
+ return compile_metal(src, sdk="macosx")
yield
diff --git a/tests/python/codegen/test_target_codegen_metal.py
b/tests/python/codegen/test_target_codegen_metal.py
index 061fe69947..fb2e7e4f38 100644
--- a/tests/python/codegen/test_target_codegen_metal.py
+++ b/tests/python/codegen/test_target_codegen_metal.py
@@ -186,7 +186,7 @@ def test_func_with_trailing_pod_params():
mod = tvm.IRModule({"main": func})
- f = tvm.compile(mod, target="metal")
+ f = tvm.tir.build(mod, target="metal")
src: str = f.imports[0].inspect_source()
occurrences = src.count("struct func_kernel_args_t")
assert occurrences == 1, occurrences
