[Mlir-commits] [mlir] [mlir][ROCDL] do not hardcode partial lld path in utilities (PR #179201)

[Oleksandr Alex Zinenko]

https://github.com/ftynse created https://github.com/llvm/llvm-project/pull/179201

`ROCDL::linkObjectCode` was unconditionally appending llvm/bin/ld.lld to the path it is been passed to to look for lld, which isn't desirable for a utility function and makes it unusable with, e.g., system lld or one from the LLVM's own build directory. Move this logic to the caller and let the utility take a full path.

>From 137cde12b16cac6b715176ba404b027287e472fb Mon Sep 17 00:00:00 2001
From: Alex Zinenko <git at ozinenko.com>
Date: Mon, 2 Feb 2026 10:54:09 +0100
Subject: [PATCH] [mlir][ROCDL] do not hardcode partial lld path in utilities

`ROCDL::linkObjectCode` was unconditionally appending llvm/bin/ld.lld to the
path it is been passed to to look for lld, which isn't desirable for a utility
function and makes it unusable with, e.g., system lld or one from the LLVM's
own build directory. Move this logic to the caller and let the utility take a
full path.
---
 mlir/include/mlir/Target/LLVM/ROCDL/Utils.h | 2 +-
 mlir/lib/Target/LLVM/ROCDL/Target.cpp       | 8 ++++----
 2 files changed, 5 insertions(+), 5 deletions(-)

diff --git a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
index 302ff6e7b8575..7f64f280cd507 100644
--- a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
+++ b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
@@ -49,7 +49,7 @@ assembleIsa(StringRef isa, StringRef targetTriple, StringRef chip,
             StringRef features, function_ref<InFlightDiagnostic()> emitError);
 
 FailureOr<SmallVector<char, 0>>
-linkObjectCode(ArrayRef<char> objectCode, StringRef toolkitPath,
+linkObjectCode(ArrayRef<char> objectCode, StringRef lldPath,
                function_ref<InFlightDiagnostic()> emitError);
 
 /// Base class for all ROCDL serializations from GPU modules into binary
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index 1af7eabfb4b10..5743acf75c9a5 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -336,7 +336,7 @@ mlir::ROCDL::assembleIsa(StringRef isa, StringRef targetTriple, StringRef chip,
 }
 
 FailureOr<SmallVector<char, 0>>
-mlir::ROCDL::linkObjectCode(ArrayRef<char> objectCode, StringRef toolkitPath,
+mlir::ROCDL::linkObjectCode(ArrayRef<char> objectCode, StringRef lldPath,
                             function_ref<InFlightDiagnostic()> emitError) {
   // Save the ISA binary to a temp file.
   int tempIsaBinaryFd = -1;
@@ -361,8 +361,6 @@ mlir::ROCDL::linkObjectCode(ArrayRef<char> objectCode, StringRef toolkitPath,
 
   llvm::FileRemover cleanupHsaco(tempHsacoFilename);
 
-  llvm::SmallString<128> lldPath(toolkitPath);
-  llvm::sys::path::append(lldPath, "llvm", "bin", "ld.lld");
   int lldResult = llvm::sys::ExecuteAndWait(
       lldPath,
       {"ld.lld", "-shared", tempIsaBinaryFilename, "-o", tempHsacoFilename});
@@ -392,8 +390,10 @@ SerializeGPUModuleBase::compileToBinary(StringRef serializedISA) {
     return failure();
 
   // Link the object code.
+  llvm::SmallString<128> lldPath(toolkitPath);
+  llvm::sys::path::append(lldPath, "llvm", "bin", "ld.lld");
   FailureOr<SmallVector<char, 0>> linkedCode =
-      ROCDL::linkObjectCode(*isaBinary, toolkitPath, errCallback);
+      ROCDL::linkObjectCode(*isaBinary, lldPath, errCallback);
   if (failed(linkedCode))
     return failure();