From ece66266d1413ca3df0db47fd5a2ee4cd42d8812 Mon Sep 17 00:00:00 2001 From: Ruihang Lai Date: Mon, 2 Feb 2026 10:13:23 -0500 Subject: [PATCH] [LLVM] Fix insertDbgValueIntrinsic for Metal backend 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 +++++++++++++++++++ .../codegen/test_gpu_codegen_allreduce.py | 5 +++- .../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 131c8212c597..ed2f927c0dd9 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 31b6511e0e66..5d73fb36bbaf 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 061fe69947e7..fb2e7e4f3888 100644 --- a/tests/python/codegen/test_target_codegen_metal.py +++ b/tests/python/codegen/test_target_codegen_metal.py @@ -186,7 +186,7 @@ def compile_metal(src, target): 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