[Mlir-commits] [mlir] [mlir][gpu] Use `SmallString`, `FailureOr` and `StringRef` in `module-to-binary` infra (PR #172284)
Ivan Butygin
llvmlistbot at llvm.org
Mon Dec 15 03:17:21 PST 2025
https://github.com/Hardcode84 created https://github.com/llvm/llvm-project/pull/172284
Instead of `std::string`, `std::optional` and `const std::string&`.
>From 26b72ccaa9b6f6b36e74757883c8bbc4373971e7 Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Mon, 15 Dec 2025 11:46:08 +0100
Subject: [PATCH 1/3] SmallString
---
mlir/include/mlir/Target/LLVM/ModuleToObject.h | 3 +--
mlir/include/mlir/Target/LLVM/ROCDL/Utils.h | 2 +-
mlir/include/mlir/Target/LLVM/XeVM/Utils.h | 2 +-
mlir/lib/Target/LLVM/ModuleToObject.cpp | 6 +++---
mlir/lib/Target/LLVM/NVVM/Target.cpp | 12 +++++-------
mlir/lib/Target/LLVM/ROCDL/Target.cpp | 4 ++--
mlir/lib/Target/LLVM/XeVM/Target.cpp | 4 ++--
7 files changed, 15 insertions(+), 18 deletions(-)
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index ea0b4b4a6ab7f..e946c1b25919c 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -45,8 +45,7 @@ class ModuleToObject {
virtual std::optional<SmallVector<char, 0>> run();
/// Translate LLVM module to textual ISA.
- /// TODO: switch to SmallString
- static FailureOr<std::string>
+ static FailureOr<SmallString<0>>
translateModuleToISA(llvm::Module &llvmModule,
llvm::TargetMachine &targetMachine,
function_ref<InFlightDiagnostic()> emitError);
diff --git a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
index 9340d3d46a9be..334aa136d3530 100644
--- a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
+++ b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
@@ -100,7 +100,7 @@ class SerializeGPUModuleBase : public LLVM::ModuleToObject {
/// Compiles assembly to a binary.
virtual std::optional<SmallVector<char, 0>>
- compileToBinary(const std::string &serializedISA);
+ compileToBinary(StringRef serializedISA);
/// Default implementation of `ModuleToObject::moduleToObject`.
std::optional<SmallVector<char, 0>>
diff --git a/mlir/include/mlir/Target/LLVM/XeVM/Utils.h b/mlir/include/mlir/Target/LLVM/XeVM/Utils.h
index 5d523f1ec457f..55a11e7f34bc8 100644
--- a/mlir/include/mlir/Target/LLVM/XeVM/Utils.h
+++ b/mlir/include/mlir/Target/LLVM/XeVM/Utils.h
@@ -40,7 +40,7 @@ class SerializeGPUModuleBase : public LLVM::ModuleToObject {
gpu::GPUModuleOp getGPUModuleOp();
/// Compiles to native code using `ocloc`.
- std::optional<SmallVector<char, 0>> compileToBinary(const std::string &asmStr,
+ std::optional<SmallVector<char, 0>> compileToBinary(StringRef asmStr,
StringRef inputFormat);
protected:
diff --git a/mlir/lib/Target/LLVM/ModuleToObject.cpp b/mlir/lib/Target/LLVM/ModuleToObject.cpp
index e20fd544ed8f7..5d779cc51d983 100644
--- a/mlir/lib/Target/LLVM/ModuleToObject.cpp
+++ b/mlir/lib/Target/LLVM/ModuleToObject.cpp
@@ -205,11 +205,11 @@ LogicalResult ModuleToObject::optimizeModule(llvm::Module &module,
return success();
}
-FailureOr<std::string> ModuleToObject::translateModuleToISA(
+FailureOr<SmallString<0>> ModuleToObject::translateModuleToISA(
llvm::Module &llvmModule, llvm::TargetMachine &targetMachine,
function_ref<InFlightDiagnostic()> emitError) {
- std::string targetISA;
- llvm::raw_string_ostream stream(targetISA);
+ SmallString<0> targetISA;
+ llvm::raw_svector_ostream stream(targetISA);
{ // Drop pstream after this to prevent the ISA from being stuck buffering
llvm::buffer_ostream pstream(stream);
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index cae5adfc086df..a5ab625b8dbb5 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -212,12 +212,10 @@ class NVPTXSerializer : public SerializeGPUModuleBase {
gpu::GPUModuleOp getOperation();
/// Compiles PTX to cubin using `ptxas`.
- std::optional<SmallVector<char, 0>>
- compileToBinary(const std::string &ptxCode);
+ std::optional<SmallVector<char, 0>> compileToBinary(StringRef ptxCode);
/// Compiles PTX to cubin using the `nvptxcompiler` library.
- std::optional<SmallVector<char, 0>>
- compileToBinaryNVPTX(const std::string &ptxCode);
+ std::optional<SmallVector<char, 0>> compileToBinaryNVPTX(StringRef ptxCode);
/// Serializes the LLVM module to an object format, depending on the
/// compilation target selected in target options.
@@ -347,7 +345,7 @@ static void setOptionalCommandlineArguments(NVVMTargetAttr target,
// TODO: clean this method & have a generic tool driver or never emit binaries
// with this mechanism and let another stage take care of it.
std::optional<SmallVector<char, 0>>
-NVPTXSerializer::compileToBinary(const std::string &ptxCode) {
+NVPTXSerializer::compileToBinary(StringRef ptxCode) {
// Determine if the serializer should create a fatbinary with the PTX embeded
// or a simple CUBIN binary.
const bool createFatbin =
@@ -570,7 +568,7 @@ NVPTXSerializer::compileToBinary(const std::string &ptxCode) {
} while (false)
std::optional<SmallVector<char, 0>>
-NVPTXSerializer::compileToBinaryNVPTX(const std::string &ptxCode) {
+NVPTXSerializer::compileToBinaryNVPTX(StringRef ptxCode) {
Location loc = getOperation().getLoc();
nvPTXCompilerHandle compiler = nullptr;
nvPTXCompileResult status;
@@ -697,7 +695,7 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
return std::nullopt;
}
moduleToObjectTimer.startTimer();
- FailureOr<std::string> serializedISA =
+ FailureOr<SmallString<0>> serializedISA =
translateModuleToISA(llvmModule, **targetMachine,
[&]() { return getOperation().emitError(); });
moduleToObjectTimer.stopTimer();
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index a1970dac8785a..e783fadb19a4e 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -381,7 +381,7 @@ mlir::ROCDL::linkObjectCode(ArrayRef<char> objectCode, StringRef toolkitPath,
}
std::optional<SmallVector<char, 0>>
-SerializeGPUModuleBase::compileToBinary(const std::string &serializedISA) {
+SerializeGPUModuleBase::compileToBinary(StringRef serializedISA) {
auto errCallback = [&]() { return getOperation().emitError(); };
// Assemble the ISA.
FailureOr<SmallVector<char, 0>> isaBinary = ROCDL::assembleIsa(
@@ -421,7 +421,7 @@ std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
}
// Translate the Module to ISA.
- FailureOr<std::string> serializedISA =
+ FailureOr<SmallString<0>> serializedISA =
translateModuleToISA(llvmModule, **targetMachine,
[&]() { return getOperation().emitError(); });
if (failed(serializedISA)) {
diff --git a/mlir/lib/Target/LLVM/XeVM/Target.cpp b/mlir/lib/Target/LLVM/XeVM/Target.cpp
index 24b518e6eed4c..221ee3b77a268 100644
--- a/mlir/lib/Target/LLVM/XeVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/XeVM/Target.cpp
@@ -108,7 +108,7 @@ gpu::GPUModuleOp SerializeGPUModuleBase::getGPUModuleOp() {
// - L0 runtime consumes IL and is external to MLIR codebase (rt wrappers).
// - `ocloc` tool can be "queried" from within MLIR.
std::optional<SmallVector<char, 0>>
-SerializeGPUModuleBase::compileToBinary(const std::string &asmStr,
+SerializeGPUModuleBase::compileToBinary(StringRef asmStr,
StringRef inputFormat) {
using TmpFile = std::pair<llvm::SmallString<128>, llvm::FileRemover>;
// Find the `ocloc` tool.
@@ -301,7 +301,7 @@ SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
// Return SPIRV if the compilation target is `assembly`.
if (targetOptions.getCompilationTarget() ==
gpu::CompilationTarget::Assembly) {
- FailureOr<std::string> serializedISA =
+ FailureOr<SmallString<0>> serializedISA =
translateModuleToISA(llvmModule, **targetMachine,
[&]() { return getGPUModuleOp().emitError(); });
if (failed(serializedISA)) {
>From ebe3adf442be48f448552027bfb29f614f1854ea Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Mon, 15 Dec 2025 11:52:07 +0100
Subject: [PATCH 2/3] FailureOr createTargetMachine()
---
.../include/mlir/Target/LLVM/ModuleToObject.h | 2 +-
mlir/lib/Target/LLVM/ModuleToObject.cpp | 34 +++++++++----------
mlir/lib/Target/LLVM/NVVM/Target.cpp | 5 ++-
mlir/lib/Target/LLVM/ROCDL/Target.cpp | 5 ++-
mlir/lib/Target/LLVM/XeVM/Target.cpp | 5 ++-
5 files changed, 23 insertions(+), 28 deletions(-)
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index e946c1b25919c..cb48c79a86487 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -81,7 +81,7 @@ class ModuleToObject {
protected:
/// Create the target machine based on the target triple and chip.
/// This can fail if the target is not available.
- std::optional<llvm::TargetMachine *> getOrCreateTargetMachine();
+ FailureOr<llvm::TargetMachine *> getOrCreateTargetMachine();
/// Loads a bitcode file from path.
std::unique_ptr<llvm::Module> loadBitcodeFile(llvm::LLVMContext &context,
diff --git a/mlir/lib/Target/LLVM/ModuleToObject.cpp b/mlir/lib/Target/LLVM/ModuleToObject.cpp
index 5d779cc51d983..a2121f527db60 100644
--- a/mlir/lib/Target/LLVM/ModuleToObject.cpp
+++ b/mlir/lib/Target/LLVM/ModuleToObject.cpp
@@ -50,8 +50,7 @@ ModuleToObject::~ModuleToObject() = default;
Operation &ModuleToObject::getOperation() { return module; }
-std::optional<llvm::TargetMachine *>
-ModuleToObject::getOrCreateTargetMachine() {
+FailureOr<llvm::TargetMachine *> ModuleToObject::getOrCreateTargetMachine() {
if (targetMachine)
return targetMachine.get();
// Load the target.
@@ -59,17 +58,17 @@ ModuleToObject::getOrCreateTargetMachine() {
llvm::Triple parsedTriple(triple);
const llvm::Target *target =
llvm::TargetRegistry::lookupTarget(parsedTriple, error);
- if (!target) {
- getOperation().emitError()
- << "Failed to lookup target for triple '" << triple << "' " << error;
- return std::nullopt;
- }
+ if (!target)
+ return getOperation().emitError()
+ << "Failed to lookup target for triple '" << triple << "' " << error;
// Create the target machine using the target.
targetMachine.reset(
target->createTargetMachine(parsedTriple, chip, features, {}, {}));
if (!targetMachine)
- return std::nullopt;
+ return getOperation().emitError()
+ << "Failed to create target machine for triple '" << triple << "'";
+
return targetMachine.get();
}
@@ -183,9 +182,8 @@ LogicalResult ModuleToObject::optimizeModule(llvm::Module &module,
return getOperation().emitError()
<< "Invalid optimization level: " << optLevel << ".";
- std::optional<llvm::TargetMachine *> targetMachine =
- getOrCreateTargetMachine();
- if (!targetMachine)
+ FailureOr<llvm::TargetMachine *> targetMachine = getOrCreateTargetMachine();
+ if (failed(targetMachine))
return getOperation().emitError()
<< "Target Machine unavailable for triple " << triple
<< ", can't optimize with LLVM\n";
@@ -226,13 +224,13 @@ FailureOr<SmallString<0>> ModuleToObject::translateModuleToISA(
void ModuleToObject::setDataLayoutAndTriple(llvm::Module &module) {
// Create the target machine.
- std::optional<llvm::TargetMachine *> targetMachine =
- getOrCreateTargetMachine();
- if (targetMachine) {
- // Set the data layout and target triple of the module.
- module.setDataLayout((*targetMachine)->createDataLayout());
- module.setTargetTriple((*targetMachine)->getTargetTriple());
- }
+ FailureOr<llvm::TargetMachine *> targetMachine = getOrCreateTargetMachine();
+ if (failed(targetMachine))
+ return;
+
+ // Set the data layout and target triple of the module.
+ module.setDataLayout((*targetMachine)->createDataLayout());
+ module.setTargetTriple((*targetMachine)->getTargetTriple());
}
std::optional<SmallVector<char, 0>>
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index a5ab625b8dbb5..0080812baf882 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -687,9 +687,8 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
#endif // LLVM_HAS_NVPTX_TARGET
// Emit PTX code.
- std::optional<llvm::TargetMachine *> targetMachine =
- getOrCreateTargetMachine();
- if (!targetMachine) {
+ FailureOr<llvm::TargetMachine *> targetMachine = getOrCreateTargetMachine();
+ if (failed(targetMachine)) {
getOperation().emitError() << "Target Machine unavailable for triple "
<< triple << ", can't optimize with LLVM\n";
return std::nullopt;
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index e783fadb19a4e..9b6f02f309f1c 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -412,9 +412,8 @@ std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
if (targetOptions.getCompilationTarget() == gpu::CompilationTarget::Offload)
return SerializeGPUModuleBase::moduleToObject(llvmModule);
- std::optional<llvm::TargetMachine *> targetMachine =
- getOrCreateTargetMachine();
- if (!targetMachine) {
+ FailureOr<llvm::TargetMachine *> targetMachine = getOrCreateTargetMachine();
+ if (failed(targetMachine)) {
getOperation().emitError() << "target Machine unavailable for triple "
<< triple << ", can't compile with LLVM";
return std::nullopt;
diff --git a/mlir/lib/Target/LLVM/XeVM/Target.cpp b/mlir/lib/Target/LLVM/XeVM/Target.cpp
index 221ee3b77a268..bd1e410db8780 100644
--- a/mlir/lib/Target/LLVM/XeVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/XeVM/Target.cpp
@@ -290,9 +290,8 @@ SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
return std::nullopt;
#endif // LLVM_HAS_SPIRV_TARGET
- std::optional<llvm::TargetMachine *> targetMachine =
- getOrCreateTargetMachine();
- if (!targetMachine) {
+ FailureOr<llvm::TargetMachine *> targetMachine = getOrCreateTargetMachine();
+ if (failed(targetMachine)) {
getGPUModuleOp().emitError() << "Target Machine unavailable for triple "
<< triple << ", can't optimize with LLVM\n";
return std::nullopt;
>From 469b62a34a95086f8606cb9f22dd06d1ec3e4e5b Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Mon, 15 Dec 2025 12:05:18 +0100
Subject: [PATCH 3/3] FailureOr
---
.../include/mlir/Target/LLVM/ModuleToObject.h | 2 +-
mlir/include/mlir/Target/LLVM/ROCDL/Utils.h | 4 +-
mlir/include/mlir/Target/LLVM/XeVM/Utils.h | 4 +-
mlir/lib/Target/LLVM/ModuleToObject.cpp | 2 +-
mlir/lib/Target/LLVM/NVVM/Target.cpp | 95 +++++++--------
mlir/lib/Target/LLVM/ROCDL/Target.cpp | 35 +++---
mlir/lib/Target/LLVM/XeVM/Target.cpp | 114 +++++++++---------
7 files changed, 120 insertions(+), 136 deletions(-)
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index cb48c79a86487..7d232dcf9ab90 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -75,7 +75,7 @@ class ModuleToObject {
/// Serializes the LLVM IR bitcode to an object file, by default it serializes
/// to LLVM bitcode.
- virtual std::optional<SmallVector<char, 0>>
+ virtual FailureOr<SmallVector<char, 0>>
moduleToObject(llvm::Module &llvmModule);
protected:
diff --git a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
index 334aa136d3530..302ff6e7b8575 100644
--- a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
+++ b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
@@ -99,11 +99,11 @@ class SerializeGPUModuleBase : public LLVM::ModuleToObject {
StringRef abiVer);
/// Compiles assembly to a binary.
- virtual std::optional<SmallVector<char, 0>>
+ virtual FailureOr<SmallVector<char, 0>>
compileToBinary(StringRef serializedISA);
/// Default implementation of `ModuleToObject::moduleToObject`.
- std::optional<SmallVector<char, 0>>
+ FailureOr<SmallVector<char, 0>>
moduleToObjectImpl(const gpu::TargetOptions &targetOptions,
llvm::Module &llvmModule);
diff --git a/mlir/include/mlir/Target/LLVM/XeVM/Utils.h b/mlir/include/mlir/Target/LLVM/XeVM/Utils.h
index 55a11e7f34bc8..455c8303a9aa8 100644
--- a/mlir/include/mlir/Target/LLVM/XeVM/Utils.h
+++ b/mlir/include/mlir/Target/LLVM/XeVM/Utils.h
@@ -40,8 +40,8 @@ class SerializeGPUModuleBase : public LLVM::ModuleToObject {
gpu::GPUModuleOp getGPUModuleOp();
/// Compiles to native code using `ocloc`.
- std::optional<SmallVector<char, 0>> compileToBinary(StringRef asmStr,
- StringRef inputFormat);
+ FailureOr<SmallVector<char, 0>> compileToBinary(StringRef asmStr,
+ StringRef inputFormat);
protected:
/// XeVM Target attribute.
diff --git a/mlir/lib/Target/LLVM/ModuleToObject.cpp b/mlir/lib/Target/LLVM/ModuleToObject.cpp
index a2121f527db60..8eaf8d9cadcd2 100644
--- a/mlir/lib/Target/LLVM/ModuleToObject.cpp
+++ b/mlir/lib/Target/LLVM/ModuleToObject.cpp
@@ -233,7 +233,7 @@ void ModuleToObject::setDataLayoutAndTriple(llvm::Module &module) {
module.setTargetTriple((*targetMachine)->getTargetTriple());
}
-std::optional<SmallVector<char, 0>>
+FailureOr<SmallVector<char, 0>>
ModuleToObject::moduleToObject(llvm::Module &llvmModule) {
SmallVector<char, 0> binaryData;
// Write the LLVM module bitcode to a buffer.
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 0080812baf882..43063d9eb5bee 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -212,14 +212,14 @@ class NVPTXSerializer : public SerializeGPUModuleBase {
gpu::GPUModuleOp getOperation();
/// Compiles PTX to cubin using `ptxas`.
- std::optional<SmallVector<char, 0>> compileToBinary(StringRef ptxCode);
+ FailureOr<SmallVector<char, 0>> compileToBinary(StringRef ptxCode);
/// Compiles PTX to cubin using the `nvptxcompiler` library.
- std::optional<SmallVector<char, 0>> compileToBinaryNVPTX(StringRef ptxCode);
+ FailureOr<SmallVector<char, 0>> compileToBinaryNVPTX(StringRef ptxCode);
/// Serializes the LLVM module to an object format, depending on the
/// compilation target selected in target options.
- std::optional<SmallVector<char, 0>>
+ FailureOr<SmallVector<char, 0>>
moduleToObject(llvm::Module &llvmModule) override;
/// Get LLVMIR->ISA performance result.
@@ -344,7 +344,7 @@ static void setOptionalCommandlineArguments(NVVMTargetAttr target,
// TODO: clean this method & have a generic tool driver or never emit binaries
// with this mechanism and let another stage take care of it.
-std::optional<SmallVector<char, 0>>
+FailureOr<SmallVector<char, 0>>
NVPTXSerializer::compileToBinary(StringRef ptxCode) {
// Determine if the serializer should create a fatbinary with the PTX embeded
// or a simple CUBIN binary.
@@ -354,12 +354,12 @@ NVPTXSerializer::compileToBinary(StringRef ptxCode) {
// Find the `ptxas` & `fatbinary` tools.
std::optional<std::string> ptxasCompiler = findTool("ptxas");
if (!ptxasCompiler)
- return std::nullopt;
+ return failure();
std::optional<std::string> fatbinaryTool;
if (createFatbin) {
fatbinaryTool = findTool("fatbinary");
if (!fatbinaryTool)
- return std::nullopt;
+ return failure();
}
Location loc = getOperation().getLoc();
@@ -371,13 +371,13 @@ NVPTXSerializer::compileToBinary(StringRef ptxCode) {
// Create temp files:
std::optional<TmpFile> ptxFile = createTemp(basename, "ptx");
if (!ptxFile)
- return std::nullopt;
+ return failure();
std::optional<TmpFile> logFile = createTemp(basename, "log");
if (!logFile)
- return std::nullopt;
+ return failure();
std::optional<TmpFile> binaryFile = createTemp(basename, "bin");
if (!binaryFile)
- return std::nullopt;
+ return failure();
TmpFile cubinFile;
if (createFatbin) {
std::string cubinFilename = (ptxFile->first + ".cubin").str();
@@ -390,17 +390,15 @@ NVPTXSerializer::compileToBinary(StringRef ptxCode) {
// Dump the PTX to a temp file.
{
llvm::raw_fd_ostream ptxStream(ptxFile->first, ec);
- if (ec) {
- emitError(loc) << "Couldn't open the file: `" << ptxFile->first
- << "`, error message: " << ec.message();
- return std::nullopt;
- }
+ if (ec)
+ return emitError(loc) << "Couldn't open the file: `" << ptxFile->first
+ << "`, error message: " << ec.message();
+
ptxStream << ptxCode;
- if (ptxStream.has_error()) {
- emitError(loc) << "An error occurred while writing the PTX to: `"
- << ptxFile->first << "`.";
- return std::nullopt;
- }
+ if (ptxStream.has_error())
+ return emitError(loc) << "An error occurred while writing the PTX to: `"
+ << ptxFile->first << "`.";
+
ptxStream.flush();
}
@@ -464,20 +462,18 @@ NVPTXSerializer::compileToBinary(StringRef ptxCode) {
// Helper function for printing tool error logs.
std::string message;
auto emitLogError =
- [&](StringRef toolName) -> std::optional<SmallVector<char, 0>> {
+ [&](StringRef toolName) -> FailureOr<SmallVector<char, 0>> {
if (message.empty()) {
llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> toolStderr =
llvm::MemoryBuffer::getFile(logFile->first);
if (toolStderr)
- emitError(loc) << toolName << " invocation failed. Log:\n"
- << toolStderr->get()->getBuffer();
+ return emitError(loc) << toolName << " invocation failed. Log:\n"
+ << toolStderr->get()->getBuffer();
else
- emitError(loc) << toolName << " invocation failed.";
- return std::nullopt;
+ return emitError(loc) << toolName << " invocation failed.";
}
- emitError(loc) << toolName
- << " invocation failed, error message: " << message;
- return std::nullopt;
+ return emitError(loc) << toolName
+ << " invocation failed, error message: " << message;
};
// Invoke PTXAS.
@@ -534,11 +530,11 @@ NVPTXSerializer::compileToBinary(StringRef ptxCode) {
// Read the fatbin.
llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> binaryBuffer =
llvm::MemoryBuffer::getFile(binaryFile->first);
- if (!binaryBuffer) {
- emitError(loc) << "Couldn't open the file: `" << binaryFile->first
- << "`, error message: " << binaryBuffer.getError().message();
- return std::nullopt;
- }
+ if (!binaryBuffer)
+ return emitError(loc) << "Couldn't open the file: `" << binaryFile->first
+ << "`, error message: "
+ << binaryBuffer.getError().message();
+
StringRef fatbin = (*binaryBuffer)->getBuffer();
return SmallVector<char, 0>(fatbin.begin(), fatbin.end());
}
@@ -567,7 +563,7 @@ NVPTXSerializer::compileToBinary(StringRef ptxCode) {
} \
} while (false)
-std::optional<SmallVector<char, 0>>
+FailureOr<SmallVector<char, 0>>
NVPTXSerializer::compileToBinaryNVPTX(StringRef ptxCode) {
Location loc = getOperation().getLoc();
nvPTXCompilerHandle compiler = nullptr;
@@ -599,13 +595,12 @@ NVPTXSerializer::compileToBinaryNVPTX(StringRef ptxCode) {
SmallVector<char> log(logSize + 1, 0);
RETURN_ON_NVPTXCOMPILER_ERROR(
nvPTXCompilerGetErrorLog(compiler, log.data()));
- emitError(loc) << "NVPTX compiler invocation failed, error log: "
- << log.data();
+ return emitError(loc)
+ << "NVPTX compiler invocation failed, error log: " << log.data();
} else {
- emitError(loc) << "NVPTX compiler invocation failed with error code: "
- << status;
+ return emitError(loc)
+ << "NVPTX compiler invocation failed with error code: " << status;
}
- return std::nullopt;
}
// Retrieve the binary.
@@ -664,7 +659,7 @@ NVPTXSerializer::compileToBinaryNVPTX(StringRef ptxCode) {
}
#endif // MLIR_ENABLE_NVPTXCOMPILER
-std::optional<SmallVector<char, 0>>
+FailureOr<SmallVector<char, 0>>
NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
llvm::Timer moduleToObjectTimer(
"moduleToObjectTimer",
@@ -681,18 +676,17 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
return SerializeGPUModuleBase::moduleToObject(llvmModule);
#if !LLVM_HAS_NVPTX_TARGET
- getOperation()->emitError(
+ return getOperation()->emitError(
"The `NVPTX` target was not built. Please enable it when building LLVM.");
- return std::nullopt;
#endif // LLVM_HAS_NVPTX_TARGET
// Emit PTX code.
FailureOr<llvm::TargetMachine *> targetMachine = getOrCreateTargetMachine();
- if (failed(targetMachine)) {
- getOperation().emitError() << "Target Machine unavailable for triple "
- << triple << ", can't optimize with LLVM\n";
- return std::nullopt;
- }
+ if (failed(targetMachine))
+ return getOperation().emitError()
+ << "Target Machine unavailable for triple " << triple
+ << ", can't optimize with LLVM\n";
+
moduleToObjectTimer.startTimer();
FailureOr<SmallString<0>> serializedISA =
translateModuleToISA(llvmModule, **targetMachine,
@@ -700,10 +694,9 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
moduleToObjectTimer.stopTimer();
llvmToISATimeInMs = moduleToObjectTimer.getTotalTime().getWallTime() * 1000;
moduleToObjectTimer.clear();
- if (failed(serializedISA)) {
- getOperation().emitError() << "Failed translating the module to ISA.";
- return std::nullopt;
- }
+ if (failed(serializedISA))
+ return getOperation().emitError()
+ << "Failed translating the module to ISA.";
if (isaCallback)
isaCallback(*serializedISA);
@@ -717,7 +710,7 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
if (targetOptions.getCompilationTarget() == gpu::CompilationTarget::Assembly)
return SmallVector<char, 0>(serializedISA->begin(), serializedISA->end());
- std::optional<SmallVector<char, 0>> result;
+ FailureOr<SmallVector<char, 0>> result;
moduleToObjectTimer.startTimer();
// Compile to binary.
#if MLIR_ENABLE_NVPTXCOMPILER
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index 9b6f02f309f1c..2669cd1ec8ca5 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -380,7 +380,7 @@ mlir::ROCDL::linkObjectCode(ArrayRef<char> objectCode, StringRef toolkitPath,
return SmallVector<char, 0>(buffer.begin(), buffer.end());
}
-std::optional<SmallVector<char, 0>>
+FailureOr<SmallVector<char, 0>>
SerializeGPUModuleBase::compileToBinary(StringRef serializedISA) {
auto errCallback = [&]() { return getOperation().emitError(); };
// Assemble the ISA.
@@ -388,18 +388,18 @@ SerializeGPUModuleBase::compileToBinary(StringRef serializedISA) {
serializedISA, this->triple, this->chip, this->features, errCallback);
if (failed(isaBinary))
- return std::nullopt;
+ return failure();
// Link the object code.
FailureOr<SmallVector<char, 0>> linkedCode =
ROCDL::linkObjectCode(*isaBinary, toolkitPath, errCallback);
if (failed(linkedCode))
- return std::nullopt;
+ return failure();
return linkedCode;
}
-std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
+FailureOr<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
const gpu::TargetOptions &targetOptions, llvm::Module &llvmModule) {
// Return LLVM IR if the compilation target is offload.
#define DEBUG_TYPE "serialize-to-llvm"
@@ -413,20 +413,18 @@ std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
return SerializeGPUModuleBase::moduleToObject(llvmModule);
FailureOr<llvm::TargetMachine *> targetMachine = getOrCreateTargetMachine();
- if (failed(targetMachine)) {
- getOperation().emitError() << "target Machine unavailable for triple "
- << triple << ", can't compile with LLVM";
- return std::nullopt;
- }
+ if (failed(targetMachine))
+ return getOperation().emitError()
+ << "target Machine unavailable for triple " << triple
+ << ", can't compile with LLVM";
// Translate the Module to ISA.
FailureOr<SmallString<0>> serializedISA =
translateModuleToISA(llvmModule, **targetMachine,
[&]() { return getOperation().emitError(); });
- if (failed(serializedISA)) {
- getOperation().emitError() << "failed translating the module to ISA";
- return std::nullopt;
- }
+ if (failed(serializedISA))
+ return getOperation().emitError() << "failed translating the module to ISA";
+
#define DEBUG_TYPE "serialize-to-isa"
LLVM_DEBUG({
llvm::dbgs() << "ISA for module: "
@@ -439,10 +437,9 @@ std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
return SmallVector<char, 0>(serializedISA->begin(), serializedISA->end());
// Compiling to binary requires a valid ROCm path, fail if it's not found.
- if (getToolkitPath().empty()) {
- getOperation().emitError() << "invalid ROCm path, please set a valid path";
- return std::nullopt;
- }
+ if (getToolkitPath().empty())
+ return getOperation().emitError()
+ << "invalid ROCm path, please set a valid path";
// Compile to binary.
return compileToBinary(*serializedISA);
@@ -455,7 +452,7 @@ class AMDGPUSerializer : public SerializeGPUModuleBase {
AMDGPUSerializer(Operation &module, ROCDLTargetAttr target,
const gpu::TargetOptions &targetOptions);
- std::optional<SmallVector<char, 0>>
+ FailureOr<SmallVector<char, 0>>
moduleToObject(llvm::Module &llvmModule) override;
private:
@@ -469,7 +466,7 @@ AMDGPUSerializer::AMDGPUSerializer(Operation &module, ROCDLTargetAttr target,
: SerializeGPUModuleBase(module, target, targetOptions),
targetOptions(targetOptions) {}
-std::optional<SmallVector<char, 0>>
+FailureOr<SmallVector<char, 0>>
AMDGPUSerializer::moduleToObject(llvm::Module &llvmModule) {
return moduleToObjectImpl(targetOptions, llvmModule);
}
diff --git a/mlir/lib/Target/LLVM/XeVM/Target.cpp b/mlir/lib/Target/LLVM/XeVM/Target.cpp
index bd1e410db8780..3ea13bdd3ea67 100644
--- a/mlir/lib/Target/LLVM/XeVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/XeVM/Target.cpp
@@ -107,51 +107,49 @@ gpu::GPUModuleOp SerializeGPUModuleBase::getGPUModuleOp() {
// There are 2 ways to access IGC: AOT (ocloc) and JIT (L0 runtime).
// - L0 runtime consumes IL and is external to MLIR codebase (rt wrappers).
// - `ocloc` tool can be "queried" from within MLIR.
-std::optional<SmallVector<char, 0>>
+FailureOr<SmallVector<char, 0>>
SerializeGPUModuleBase::compileToBinary(StringRef asmStr,
StringRef inputFormat) {
using TmpFile = std::pair<llvm::SmallString<128>, llvm::FileRemover>;
// Find the `ocloc` tool.
std::optional<std::string> oclocCompiler = findTool("ocloc");
if (!oclocCompiler)
- return std::nullopt;
+ return failure();
Location loc = getGPUModuleOp().getLoc();
std::string basename = llvm::formatv(
"mlir-{0}-{1}-{2}", getGPUModuleOp().getNameAttr().getValue(),
getTarget().getTriple(), getTarget().getChip());
auto createTemp = [&](StringRef name,
- StringRef suffix) -> std::optional<TmpFile> {
+ StringRef suffix) -> FailureOr<TmpFile> {
llvm::SmallString<128> filePath;
- if (auto ec = llvm::sys::fs::createTemporaryFile(name, suffix, filePath)) {
- getGPUModuleOp().emitError()
- << "Couldn't create the temp file: `" << filePath
- << "`, error message: " << ec.message();
- return std::nullopt;
- }
+ if (auto ec = llvm::sys::fs::createTemporaryFile(name, suffix, filePath))
+ return getGPUModuleOp().emitError()
+ << "Couldn't create the temp file: `" << filePath
+ << "`, error message: " << ec.message();
+
return TmpFile(filePath, llvm::FileRemover(filePath.c_str()));
};
// Create temp file
- std::optional<TmpFile> asmFile = createTemp(basename, "asm");
- std::optional<TmpFile> binFile = createTemp(basename, "");
- std::optional<TmpFile> logFile = createTemp(basename, "log");
- if (!logFile || !asmFile || !binFile)
- return std::nullopt;
+ FailureOr<TmpFile> asmFile = createTemp(basename, "asm");
+ FailureOr<TmpFile> binFile = createTemp(basename, "");
+ FailureOr<TmpFile> logFile = createTemp(basename, "log");
+ if (failed(logFile) || failed(asmFile) || failed(binFile))
+ return failure();
// Dump the assembly to a temp file
std::error_code ec;
{
llvm::raw_fd_ostream asmStream(asmFile->first, ec);
- if (ec) {
- emitError(loc) << "Couldn't open the file: `" << asmFile->first
- << "`, error message: " << ec.message();
- return std::nullopt;
- }
+ if (ec)
+ return emitError(loc) << "Couldn't open the file: `" << asmFile->first
+ << "`, error message: " << ec.message();
+
asmStream << asmStr;
- if (asmStream.has_error()) {
- emitError(loc) << "An error occurred while writing the assembly to: `"
- << asmFile->first << "`.";
- return std::nullopt;
- }
+ if (asmStream.has_error())
+ return emitError(loc)
+ << "An error occurred while writing the assembly to: `"
+ << asmFile->first << "`.";
+
asmStream.flush();
}
// Set cmd options
@@ -176,20 +174,18 @@ SerializeGPUModuleBase::compileToBinary(StringRef asmStr,
// Helper function for printing tool error logs.
std::string message;
auto emitLogError =
- [&](StringRef toolName) -> std::optional<SmallVector<char, 0>> {
+ [&](StringRef toolName) -> FailureOr<SmallVector<char, 0>> {
if (message.empty()) {
llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> toolStderr =
llvm::MemoryBuffer::getFile(logFile->first);
if (toolStderr)
- emitError(loc) << toolName << " invocation failed. Log:\n"
- << toolStderr->get()->getBuffer();
+ return emitError(loc) << toolName << " invocation failed. Log:\n"
+ << toolStderr->get()->getBuffer();
else
- emitError(loc) << toolName << " invocation failed.";
- return std::nullopt;
+ return emitError(loc) << toolName << " invocation failed.";
}
- emitError(loc) << toolName
- << " invocation failed, error message: " << message;
- return std::nullopt;
+ return emitError(loc) << toolName
+ << " invocation failed, error message: " << message;
};
std::optional<StringRef> redirects[] = {
std::nullopt,
@@ -203,11 +199,11 @@ SerializeGPUModuleBase::compileToBinary(StringRef asmStr,
binFile->first.append(".bin");
llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> binaryBuffer =
llvm::MemoryBuffer::getFile(binFile->first);
- if (!binaryBuffer) {
- emitError(loc) << "Couldn't open the file: `" << binFile->first
- << "`, error message: " << binaryBuffer.getError().message();
- return std::nullopt;
- }
+ if (!binaryBuffer)
+ return emitError(loc) << "Couldn't open the file: `" << binFile->first
+ << "`, error message: "
+ << binaryBuffer.getError().message();
+
StringRef bin = (*binaryBuffer)->getBuffer();
return SmallVector<char, 0>(bin.begin(), bin.end());
}
@@ -245,7 +241,7 @@ class SPIRVSerializer : public SerializeGPUModuleBase {
/// Serializes the LLVM module to an object format, depending on the
/// compilation target selected in target options.
- std::optional<SmallVector<char, 0>>
+ FailureOr<SmallVector<char, 0>>
moduleToObject(llvm::Module &llvmModule) override;
private:
@@ -269,7 +265,7 @@ void SPIRVSerializer::init() {
});
}
-std::optional<SmallVector<char, 0>>
+FailureOr<SmallVector<char, 0>>
SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
#define DEBUG_TYPE "serialize-to-llvm"
LLVM_DEBUG({
@@ -285,17 +281,16 @@ SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
return SerializeGPUModuleBase::moduleToObject(llvmModule);
#if !LLVM_HAS_SPIRV_TARGET
- getGPUModuleOp()->emitError("The `SPIRV` target was not built. Please enable "
- "it when building LLVM.");
- return std::nullopt;
+ return getGPUModuleOp()->emitError(
+ "The `SPIRV` target was not built. Please enable "
+ "it when building LLVM.");
#endif // LLVM_HAS_SPIRV_TARGET
FailureOr<llvm::TargetMachine *> targetMachine = getOrCreateTargetMachine();
- if (failed(targetMachine)) {
- getGPUModuleOp().emitError() << "Target Machine unavailable for triple "
- << triple << ", can't optimize with LLVM\n";
- return std::nullopt;
- }
+ if (failed(targetMachine))
+ return getGPUModuleOp().emitError()
+ << "Target Machine unavailable for triple " << triple
+ << ", can't optimize with LLVM\n";
// Return SPIRV if the compilation target is `assembly`.
if (targetOptions.getCompilationTarget() ==
@@ -303,11 +298,10 @@ SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
FailureOr<SmallString<0>> serializedISA =
translateModuleToISA(llvmModule, **targetMachine,
[&]() { return getGPUModuleOp().emitError(); });
- if (failed(serializedISA)) {
- getGPUModuleOp().emitError() << "Failed translating the module to ISA."
- << triple << ", can't compile with LLVM\n";
- return std::nullopt;
- }
+ if (failed(serializedISA))
+ return getGPUModuleOp().emitError()
+ << "Failed translating the module to ISA." << triple
+ << ", can't compile with LLVM\n";
#define DEBUG_TYPE "serialize-to-isa"
LLVM_DEBUG({
@@ -330,14 +324,14 @@ SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
// implementation switches to native XeVM binary format.
std::optional<std::string> serializedSPIRVBinary =
translateToSPIRVBinary(llvmModule, **targetMachine);
- if (!serializedSPIRVBinary) {
- getGPUModuleOp().emitError() << "Failed translating the module to Binary.";
- return std::nullopt;
- }
- if (serializedSPIRVBinary->size() % 4) {
- getGPUModuleOp().emitError() << "SPIRV code size must be a multiple of 4.";
- return std::nullopt;
- }
+ if (!serializedSPIRVBinary)
+ return getGPUModuleOp().emitError()
+ << "Failed translating the module to Binary.";
+
+ if (serializedSPIRVBinary->size() % 4)
+ return getGPUModuleOp().emitError()
+ << "SPIRV code size must be a multiple of 4.";
+
StringRef bin(serializedSPIRVBinary->c_str(), serializedSPIRVBinary->size());
return SmallVector<char, 0>(bin.begin(), bin.end());
}
More information about the Mlir-commits
mailing list