[LLVM] Fix insertDbgValueIntrinsic for Metal backend

[MasterJH5574]

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.

[gemini-code-assist]

Summary of Changes

Hello @MasterJH5574, I'm Gemini Code Assist¹! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!

This pull request resolves a critical compilation issue affecting the Metal backend when using specific LLVM versions. It refines how debug information is handled for different data types, ensuring compatibility with LLVM's stricter requirements for debug intrinsics. The changes enable successful compilation of previously failing TIR code for Metal.

Highlights

• LLVM Debug Intrinsic Fix: Implemented a conditional fix in CodeGenLLVM::AddDebugInformation to correctly use insertDbgValueIntrinsic for non-pointer types (floats, vectors) when compiling with LLVM 15 or newer, addressing a Metal backend compilation issue.
• LLVM Version Compatibility: Adapted the insertDbgValueIntrinsic call signature for LLVM 20+ to use BasicBlock::iterator as the insertion point, ensuring forward compatibility.
• Test Updates: Updated test_gpu_codegen_allreduce.py to include tvm_ffi import and refactor compile_metal import, and test_target_codegen_metal.py to use tvm.tir.build instead of tvm.compile.

🧠 New Feature in Public Preview: You can now enable Memory to help Gemini Code Assist learn from your team's feedback. This makes future code reviews more consistent and personalized to your project's style. Click here to enable Memory in your admin console.

Using Gemini Code Assist

The full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips.

Invoking Gemini

You can request assistance from Gemini at any point by creating a comment using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands on the current page.

Feature	Command	Description
Code Review	/gemini review	Performs a code review for the current pull request in its current state.
Pull Request Summary	/gemini summary	Provides a summary of the current pull request in its current state.
Comment	@gemini-code-assist	Responds in comments when explicitly tagged, both in pull request comments and review comments.
Help	/gemini help	Displays a list of available commands.

Customization

To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a .gemini/ folder in the base of the repository. Detailed instructions can be found here.

Limitations & Feedback

Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with 👍 and 👎 on @gemini-code-assist comments. If you're interested in giving your feedback about your experience with Gemini Code Assist for Github and other Google products, sign up here.

You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the Gemini Code Assist IDE Extension.

Footnotes

1. Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configure Gemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution. ↩

[gemini-code-assist]

Code Review

This pull request aims to fix an LLVM issue for the Metal backend related to debug information generation. The main change in codegen_llvm.cc introduces the use of insertDbgValueIntrinsic for non-pointer types on LLVM 15+, which is correct in principle. However, the current implementation has two critical issues: it removes the existing insertDeclare logic, causing a regression for pointer types and older LLVM versions, and it has a bug in the new logic for LLVM versions 15-19. I've provided a detailed comment with a suggested fix. The changes in the Python test files seem fine.

[MasterJH5574]

/gemini review

[gemini-code-assist]

Code Review

This pull request addresses a compilation issue for the Metal backend on LLVM 15+ by using insertDbgValueIntrinsic for non-pointer types in debug information generation. This change correctly handles SSA values that don't have a memory address, resolving an LLVM assertion failure. The changes in src/target/llvm/codegen_llvm.cc are well-structured with preprocessor directives to maintain compatibility across different LLVM versions. Additionally, unit tests have been updated to reflect new API usage and to include compile-only tests for Metal to verify the fix. My review includes a minor suggestion to improve the accuracy of a comment. Overall, the changes look good and effectively fix the described issue.