diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp b/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
index be04a42..88ee535 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
@@ -221,7 +221,12 @@
     // Serialize hsaco kernel into the binary that we will embed in the
     // final FlatBuffer.
     std::string targetObj = translateModuleToObj(*llvmModule, *targetMachine);
-    std::string targetHSACO = createHsaco(targetObj, libraryName);
+    std::string targetHSACO =
+        createHsaco(variantOp.getLoc(), targetObj, libraryName);
+    if (targetHSACO.empty()) {
+      return failure();
+    }
+
     if (!options.dumpBinariesPath.empty()) {
       dumpDataToPath(options.dumpBinariesPath, options.dumpBaseName,
                      variantOp.getName(), ".hsaco", targetHSACO);
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.h b/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.h
index 7077ac0..04eca04 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.h
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.h
@@ -23,7 +23,7 @@
                           std::string bitCodeDir);
 
 // Compiles ISAToHsaco Code
-std::string createHsaco(const std::string isa, StringRef name);
+std::string createHsaco(Location loc, const std::string isa, StringRef name);
 
 } // namespace HAL
 } // namespace IREE
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTargetUtils.cpp b/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTargetUtils.cpp
index c3064c9..da6814d 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTargetUtils.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTargetUtils.cpp
@@ -5,6 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
 #include "iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.h"
+#include "iree/compiler/Utils/ToolUtils.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IRReader/IRReader.h"
 #include "llvm/Linker/Linker.h"
@@ -131,15 +132,14 @@
 // Inspiration from this section comes from LLVM-PROJECT-MLIR by
 // ROCmSoftwarePlatform
 // https://github.com/ROCmSoftwarePlatform/llvm-project-mlir/blob/miopen-dialect/mlir/lib/ExecutionEngine/ROCm/BackendUtils.cpp
-std::string createHsaco(const std::string isa, StringRef name) {
+std::string createHsaco(Location loc, const std::string isa, StringRef name) {
   // Save the ISA binary to a temp file.
   int tempIsaBinaryFd = -1;
   SmallString<128> tempIsaBinaryFilename;
   std::error_code ec = llvm::sys::fs::createTemporaryFile(
       "kernel", "o", tempIsaBinaryFd, tempIsaBinaryFilename);
   if (ec) {
-    llvm::WithColor::error(llvm::errs(), name)
-        << "temporary file for ISA binary creation error.\n";
+    mlir::emitError(loc) << "temporary file for ISA binary creation error";
     return {};
   }
   llvm::FileRemover cleanupIsaBinary(tempIsaBinaryFilename);
@@ -153,56 +153,45 @@
   ec = llvm::sys::fs::createTemporaryFile("kernel", "hsaco", tempHsacoFD,
                                           tempHsacoFilename);
   if (ec) {
-    llvm::WithColor::error(llvm::errs(), name)
-        << "temporary file for HSA code object creation error.\n";
+    mlir::emitError(loc) << "temporary file for HSA code object creation error";
     return {};
   }
   llvm::FileRemover cleanupHsaco(tempHsacoFilename);
 
   // Invoke lld. Expect a true return value from lld.
   // Searching for LLD
-  std::string lldProgram;
-  std::string toolName = "ld.lld";
-  if (llvm::sys::fs::exists(toolName)) {
-    llvm::SmallString<256> absolutePath(toolName);
-    llvm::sys::fs::make_absolute(absolutePath);
-    lldProgram = std::string(absolutePath);
-  } else {
-    // Next search the environment path.
-    if (auto result = llvm::sys::Process::FindInEnvPath("PATH", toolName)) {
-      lldProgram = std::string(*result);
-    }
-  }
+  const SmallVector<std::string> &toolNames{"iree-lld", "lld"};
+  std::string lldProgram = findTool(toolNames);
   if (lldProgram.empty()) {
-    llvm::WithColor::error(llvm::errs(), name)
-        << "unable to find ld.lld in PATH\n";
+    mlir::emitError(loc) << "unable to find iree-lld";
     return {};
   }
-  // Setting Up LLD Args
   std::vector<llvm::StringRef> lldArgs{
-      llvm::StringRef("ld.lld"),   llvm::StringRef("-flavor"),
-      llvm::StringRef("gnu"),      llvm::StringRef("-shared"),
-      tempIsaBinaryFilename.str(), llvm::StringRef("-o"),
+      lldProgram,
+      llvm::StringRef("-flavor"),
+      llvm::StringRef("gnu"),
+      llvm::StringRef("-shared"),
+      tempIsaBinaryFilename.str(),
+      llvm::StringRef("-o"),
       tempHsacoFilename.str(),
   };
 
   // Executing LLD
   std::string errorMessage;
   int lldResult = llvm::sys::ExecuteAndWait(
-      lldProgram, llvm::ArrayRef<llvm::StringRef>(lldArgs), std::nullopt, {}, 5,
-      0, &errorMessage);
+      unescapeCommandLineComponent(lldProgram),
+      llvm::ArrayRef<llvm::StringRef>(lldArgs),
+      llvm::StringRef("LLD_VERSION=IREE"), {}, 0, 0, &errorMessage);
   if (lldResult) {
-    llvm::WithColor::error(llvm::errs(), name)
-        << "ld.lld execute fail:" << errorMessage << "Error Code:" << lldResult
-        << "\n";
+    mlir::emitError(loc) << "iree-lld execute fail:" << errorMessage
+                         << "Error Code:" << lldResult;
     return {};
   }
 
   // Load the HSA code object.
   auto hsacoFile = mlir::openInputFile(tempHsacoFilename);
   if (!hsacoFile) {
-    llvm::WithColor::error(llvm::errs(), name)
-        << "read HSA code object from temp file error.\n";
+    mlir::emitError(loc) << "read HSA code object from temp file error";
     return {};
   }
   std::string strHSACO(hsacoFile->getBuffer().begin(),