[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
#18585 (comment).

Unit tests are updated, though they are not enabled in CI for now.

Author: MasterJH5574

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(

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
 

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