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

[llvmlistbot at llvm.org]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir

Author: Oleksandr "Alex" Zinenko (ftynse)

<details>
<summary>Changes</summary>

`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.

---
Full diff: https://github.com/llvm/llvm-project/pull/179201.diff


2 Files Affected:

- (modified) mlir/include/mlir/Target/LLVM/ROCDL/Utils.h (+1-1) 
- (modified) mlir/lib/Target/LLVM/ROCDL/Target.cpp (+4-4) 


``````````diff
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();
 

``````````

</details>


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