[Mlir-commits] [mlir] [mlir][gpu] Expose some utility functions from `gpu-to-binary` infra (PR #172205)
Ivan Butygin
llvmlistbot at llvm.org
Sun Dec 14 04:47:00 PST 2025
https://github.com/Hardcode84 updated https://github.com/llvm/llvm-project/pull/172205
>From daaccb6bbb9f440497281bfa9bf39cf25d7d5f98 Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Sun, 14 Dec 2025 12:05:35 +0100
Subject: [PATCH 1/7] translateToISA
---
mlir/include/mlir/Target/LLVM/ModuleToObject.h | 10 +++++-----
mlir/lib/Target/LLVM/ModuleToObject.cpp | 8 ++++----
mlir/lib/Target/LLVM/NVVM/Target.cpp | 8 ++++----
mlir/lib/Target/LLVM/ROCDL/Target.cpp | 6 +++---
mlir/lib/Target/LLVM/XeVM/Target.cpp | 6 +++---
5 files changed, 19 insertions(+), 19 deletions(-)
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index 11fea6f0a4443..a8552272a7723 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -23,6 +23,11 @@ class TargetMachine;
namespace mlir {
namespace LLVM {
+
+/// Utility function for translating to ISA.
+FailureOr<std::string> translateModuleToISA(llvm::Module &llvmModule,
+ llvm::TargetMachine &targetMachine);
+
class ModuleTranslation;
/// Utility base class for transforming operations into binary objects, by
/// default it returns the serialized LLVM bitcode for the module. The
@@ -98,11 +103,6 @@ class ModuleToObject {
/// Optimize the module.
virtual LogicalResult optimizeModule(llvm::Module &module, int optL);
- /// Utility function for translating to ISA, returns `std::nullopt` on
- /// failure.
- static std::optional<std::string>
- translateToISA(llvm::Module &llvmModule, llvm::TargetMachine &targetMachine);
-
protected:
/// Module to transform to a binary object.
Operation &module;
diff --git a/mlir/lib/Target/LLVM/ModuleToObject.cpp b/mlir/lib/Target/LLVM/ModuleToObject.cpp
index 4098ccc548dc1..221184d2fcdd4 100644
--- a/mlir/lib/Target/LLVM/ModuleToObject.cpp
+++ b/mlir/lib/Target/LLVM/ModuleToObject.cpp
@@ -205,9 +205,9 @@ LogicalResult ModuleToObject::optimizeModule(llvm::Module &module,
return success();
}
-std::optional<std::string>
-ModuleToObject::translateToISA(llvm::Module &llvmModule,
- llvm::TargetMachine &targetMachine) {
+FailureOr<std::string>
+mlir::LLVM::translateModuleToISA(llvm::Module &llvmModule,
+ llvm::TargetMachine &targetMachine) {
std::string targetISA;
llvm::raw_string_ostream stream(targetISA);
@@ -217,7 +217,7 @@ ModuleToObject::translateToISA(llvm::Module &llvmModule,
if (targetMachine.addPassesToEmitFile(codegenPasses, pstream, nullptr,
llvm::CodeGenFileType::AssemblyFile))
- return std::nullopt;
+ return failure();
codegenPasses.run(llvmModule);
}
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 8760ea8588e2c..be68f1f8f42b4 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -697,18 +697,18 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
return std::nullopt;
}
moduleToObjectTimer.startTimer();
- std::optional<std::string> serializedISA =
- translateToISA(llvmModule, **targetMachine);
+ FailureOr<std::string> serializedISA =
+ mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine);
moduleToObjectTimer.stopTimer();
llvmToISATimeInMs = moduleToObjectTimer.getTotalTime().getWallTime() * 1000;
moduleToObjectTimer.clear();
- if (!serializedISA) {
+ if (failed(serializedISA)) {
getOperation().emitError() << "Failed translating the module to ISA.";
return std::nullopt;
}
if (isaCallback)
- isaCallback(serializedISA.value());
+ isaCallback(*serializedISA);
#define DEBUG_TYPE "serialize-to-isa"
LDBG() << "PTX for module: " << getOperation().getNameAttr() << "\n"
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index f813f8db8fc94..5c0b8da2037ab 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -422,9 +422,9 @@ std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
}
// Translate the Module to ISA.
- std::optional<std::string> serializedISA =
- translateToISA(llvmModule, **targetMachine);
- if (!serializedISA) {
+ FailureOr<std::string> serializedISA =
+ mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine);
+ if (failed(serializedISA)) {
getOperation().emitError() << "failed translating the module to ISA";
return std::nullopt;
}
diff --git a/mlir/lib/Target/LLVM/XeVM/Target.cpp b/mlir/lib/Target/LLVM/XeVM/Target.cpp
index 1e6784a25d64a..d78475cbd0c22 100644
--- a/mlir/lib/Target/LLVM/XeVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/XeVM/Target.cpp
@@ -301,9 +301,9 @@ SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
// Return SPIRV if the compilation target is `assembly`.
if (targetOptions.getCompilationTarget() ==
gpu::CompilationTarget::Assembly) {
- std::optional<std::string> serializedISA =
- translateToISA(llvmModule, **targetMachine);
- if (!serializedISA) {
+ FailureOr<std::string> serializedISA =
+ mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine);
+ if (failed(serializedISA)) {
getGPUModuleOp().emitError() << "Failed translating the module to ISA."
<< triple << ", can't compile with LLVM\n";
return std::nullopt;
>From 6722f7e81f5ad0bbf8bf6b96e2a0685d2b8dcb16 Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Sun, 14 Dec 2025 12:29:21 +0100
Subject: [PATCH 2/7] asseble ISA
---
mlir/include/mlir/Target/LLVM/ROCDL/Utils.h | 8 +++---
mlir/lib/Target/LLVM/ROCDL/Target.cpp | 28 +++++++++------------
2 files changed, 17 insertions(+), 19 deletions(-)
diff --git a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
index 8f5d4162984fa..aa97ecf944f46 100644
--- a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
+++ b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
@@ -43,6 +43,11 @@ enum class AMDGCNLibraries : uint32_t {
All = (LastLib << 1) - 1
};
+/// Assembles ISA to an object code.
+FailureOr<SmallVector<char, 0>>
+assembleIsa(StringRef isa, StringRef targetTriple, StringRef chip,
+ StringRef features, function_ref<InFlightDiagnostic()> emitError);
+
/// Base class for all ROCDL serializations from GPU modules into binary
/// strings. By default this class serializes into LLVM bitcode.
class SerializeGPUModuleBase : public LLVM::ModuleToObject {
@@ -98,9 +103,6 @@ class SerializeGPUModuleBase : public LLVM::ModuleToObject {
moduleToObjectImpl(const gpu::TargetOptions &targetOptions,
llvm::Module &llvmModule);
- /// Returns the assembled ISA.
- std::optional<SmallVector<char, 0>> assembleIsa(StringRef isa);
-
/// ROCDL target attribute.
ROCDLTargetAttr target;
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index 5c0b8da2037ab..bb77712aa6519 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -277,12 +277,10 @@ void SerializeGPUModuleBase::addControlVariables(
}
}
-std::optional<SmallVector<char, 0>>
-SerializeGPUModuleBase::assembleIsa(StringRef isa) {
- auto loc = getOperation().getLoc();
-
- StringRef targetTriple = this->triple;
-
+FailureOr<SmallVector<char, 0>>
+mlir::ROCDL::assembleIsa(StringRef isa, StringRef targetTriple, StringRef chip,
+ StringRef features,
+ function_ref<InFlightDiagnostic()> emitError) {
SmallVector<char, 0> result;
llvm::raw_svector_ostream os(result);
@@ -290,10 +288,8 @@ SerializeGPUModuleBase::assembleIsa(StringRef isa) {
std::string error;
const llvm::Target *target =
llvm::TargetRegistry::lookupTarget(triple, error);
- if (!target) {
- emitError(loc, Twine("failed to lookup target: ") + error);
- return std::nullopt;
- }
+ if (!target)
+ return emitError() << "failed to lookup target: " << error;
llvm::SourceMgr srcMgr;
srcMgr.AddNewSourceBuffer(llvm::MemoryBuffer::getMemBuffer(isa), SMLoc());
@@ -330,10 +326,8 @@ SerializeGPUModuleBase::assembleIsa(StringRef isa) {
std::unique_ptr<llvm::MCTargetAsmParser> tap(
target->createMCAsmParser(*sti, *parser, *mcii, mcOptions));
- if (!tap) {
- emitError(loc, "assembler initialization error");
- return std::nullopt;
- }
+ if (!tap)
+ return emitError() << "assembler initialization error";
parser->setTargetParser(*tap);
parser->Run(false);
@@ -343,9 +337,11 @@ SerializeGPUModuleBase::assembleIsa(StringRef isa) {
std::optional<SmallVector<char, 0>>
SerializeGPUModuleBase::compileToBinary(const std::string &serializedISA) {
// Assemble the ISA.
- std::optional<SmallVector<char, 0>> isaBinary = assembleIsa(serializedISA);
+ FailureOr<SmallVector<char, 0>> isaBinary = ROCDL::assembleIsa(
+ serializedISA, this->triple, this->chip, this->features,
+ [&]() { return getOperation().emitError(); });
- if (!isaBinary) {
+ if (failed(isaBinary)) {
getOperation().emitError() << "failed during ISA assembling";
return std::nullopt;
}
>From 1b02c5fa1f614d69a89e54d89e169c4b9c5406f1 Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Sun, 14 Dec 2025 12:40:26 +0100
Subject: [PATCH 3/7] link
---
mlir/include/mlir/Target/LLVM/ROCDL/Utils.h | 4 ++
mlir/lib/Target/LLVM/ROCDL/Target.cpp | 69 +++++++++++----------
2 files changed, 40 insertions(+), 33 deletions(-)
diff --git a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
index aa97ecf944f46..9340d3d46a9be 100644
--- a/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
+++ b/mlir/include/mlir/Target/LLVM/ROCDL/Utils.h
@@ -48,6 +48,10 @@ FailureOr<SmallVector<char, 0>>
assembleIsa(StringRef isa, StringRef targetTriple, StringRef chip,
StringRef features, function_ref<InFlightDiagnostic()> emitError);
+FailureOr<SmallVector<char, 0>>
+linkObjectCode(ArrayRef<char> objectCode, StringRef toolkitPath,
+ function_ref<InFlightDiagnostic()> emitError);
+
/// Base class for all ROCDL serializations from GPU modules into binary
/// strings. By default this class serializes into LLVM bitcode.
class SerializeGPUModuleBase : public LLVM::ModuleToObject {
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index bb77712aa6519..f9269e0934271 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -334,42 +334,30 @@ mlir::ROCDL::assembleIsa(StringRef isa, StringRef targetTriple, StringRef chip,
return std::move(result);
}
-std::optional<SmallVector<char, 0>>
-SerializeGPUModuleBase::compileToBinary(const std::string &serializedISA) {
- // Assemble the ISA.
- FailureOr<SmallVector<char, 0>> isaBinary = ROCDL::assembleIsa(
- serializedISA, this->triple, this->chip, this->features,
- [&]() { return getOperation().emitError(); });
-
- if (failed(isaBinary)) {
- getOperation().emitError() << "failed during ISA assembling";
- return std::nullopt;
- }
-
+FailureOr<SmallVector<char, 0>>
+mlir::ROCDL::linkObjectCode(ArrayRef<char> objectCode, StringRef toolkitPath,
+ function_ref<InFlightDiagnostic()> emitError) {
// Save the ISA binary to a temp file.
int tempIsaBinaryFd = -1;
SmallString<128> tempIsaBinaryFilename;
if (llvm::sys::fs::createTemporaryFile("kernel%%", "o", tempIsaBinaryFd,
- tempIsaBinaryFilename)) {
- getOperation().emitError()
- << "failed to create a temporary file for dumping the ISA binary";
- return std::nullopt;
- }
+ tempIsaBinaryFilename))
+ return emitError()
+ << "failed to create a temporary file for dumping the ISA binary";
+
llvm::FileRemover cleanupIsaBinary(tempIsaBinaryFilename);
{
llvm::raw_fd_ostream tempIsaBinaryOs(tempIsaBinaryFd, true);
- tempIsaBinaryOs << StringRef(isaBinary->data(), isaBinary->size());
+ tempIsaBinaryOs << StringRef(objectCode.data(), objectCode.size());
tempIsaBinaryOs.flush();
}
// Create a temp file for HSA code object.
SmallString<128> tempHsacoFilename;
- if (llvm::sys::fs::createTemporaryFile("kernel", "hsaco",
- tempHsacoFilename)) {
- getOperation().emitError()
- << "failed to create a temporary file for the HSA code object";
- return std::nullopt;
- }
+ if (llvm::sys::fs::createTemporaryFile("kernel", "hsaco", tempHsacoFilename))
+ return emitError()
+ << "failed to create a temporary file for the HSA code object";
+
llvm::FileRemover cleanupHsaco(tempHsacoFilename);
llvm::SmallString<128> lldPath(toolkitPath);
@@ -377,25 +365,40 @@ SerializeGPUModuleBase::compileToBinary(const std::string &serializedISA) {
int lldResult = llvm::sys::ExecuteAndWait(
lldPath,
{"ld.lld", "-shared", tempIsaBinaryFilename, "-o", tempHsacoFilename});
- if (lldResult != 0) {
- getOperation().emitError() << "lld invocation failed";
- return std::nullopt;
- }
+ if (lldResult != 0)
+ return emitError() << "lld invocation failed";
// Load the HSA code object.
auto hsacoFile =
llvm::MemoryBuffer::getFile(tempHsacoFilename, /*IsText=*/false);
- if (!hsacoFile) {
- getOperation().emitError()
- << "failed to read the HSA code object from the temp file";
- return std::nullopt;
- }
+ if (!hsacoFile)
+ return emitError()
+ << "failed to read the HSA code object from the temp file";
StringRef buffer = (*hsacoFile)->getBuffer();
return SmallVector<char, 0>(buffer.begin(), buffer.end());
}
+std::optional<SmallVector<char, 0>>
+SerializeGPUModuleBase::compileToBinary(const std::string &serializedISA) {
+ auto errCallback = [&]() { return getOperation().emitError(); };
+ // Assemble the ISA.
+ FailureOr<SmallVector<char, 0>> isaBinary = ROCDL::assembleIsa(
+ serializedISA, this->triple, this->chip, this->features, errCallback);
+
+ if (failed(isaBinary))
+ return std::nullopt;
+
+ // Link the object code.
+ FailureOr<SmallVector<char, 0>> linkedCode =
+ ROCDL::linkObjectCode(*isaBinary, toolkitPath, errCallback);
+ if (failed(linkedCode))
+ return std::nullopt;
+
+ return linkedCode;
+}
+
std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
const gpu::TargetOptions &targetOptions, llvm::Module &llvmModule) {
// Return LLVM IR if the compilation target is offload.
>From 8e20bc0281f5a20e2b76922634bbb177589cc64c Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Sun, 14 Dec 2025 12:43:51 +0100
Subject: [PATCH 4/7] err handler
---
mlir/include/mlir/Target/LLVM/ModuleToObject.h | 6 ++++--
mlir/lib/Target/LLVM/ModuleToObject.cpp | 5 +++--
mlir/lib/Target/LLVM/NVVM/Target.cpp | 4 +++-
mlir/lib/Target/LLVM/ROCDL/Target.cpp | 4 +++-
mlir/lib/Target/LLVM/XeVM/Target.cpp | 4 +++-
5 files changed, 16 insertions(+), 7 deletions(-)
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index a8552272a7723..5be399065ef74 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -25,8 +25,10 @@ namespace mlir {
namespace LLVM {
/// Utility function for translating to ISA.
-FailureOr<std::string> translateModuleToISA(llvm::Module &llvmModule,
- llvm::TargetMachine &targetMachine);
+FailureOr<std::string>
+translateModuleToISA(llvm::Module &llvmModule,
+ llvm::TargetMachine &targetMachine,
+ function_ref<InFlightDiagnostic()> emitError);
class ModuleTranslation;
/// Utility base class for transforming operations into binary objects, by
diff --git a/mlir/lib/Target/LLVM/ModuleToObject.cpp b/mlir/lib/Target/LLVM/ModuleToObject.cpp
index 221184d2fcdd4..eb99e69b626b6 100644
--- a/mlir/lib/Target/LLVM/ModuleToObject.cpp
+++ b/mlir/lib/Target/LLVM/ModuleToObject.cpp
@@ -207,7 +207,8 @@ LogicalResult ModuleToObject::optimizeModule(llvm::Module &module,
FailureOr<std::string>
mlir::LLVM::translateModuleToISA(llvm::Module &llvmModule,
- llvm::TargetMachine &targetMachine) {
+ llvm::TargetMachine &targetMachine,
+ function_ref<InFlightDiagnostic()> emitError) {
std::string targetISA;
llvm::raw_string_ostream stream(targetISA);
@@ -217,7 +218,7 @@ mlir::LLVM::translateModuleToISA(llvm::Module &llvmModule,
if (targetMachine.addPassesToEmitFile(codegenPasses, pstream, nullptr,
llvm::CodeGenFileType::AssemblyFile))
- return failure();
+ return emitError() << "Target machine cannot emit assembly";
codegenPasses.run(llvmModule);
}
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index be68f1f8f42b4..c22f8a23d2b2b 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -698,7 +698,9 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
}
moduleToObjectTimer.startTimer();
FailureOr<std::string> serializedISA =
- mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine);
+ mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine, [&]() {
+ return getOperation().emitError();
+ });
moduleToObjectTimer.stopTimer();
llvmToISATimeInMs = moduleToObjectTimer.getTotalTime().getWallTime() * 1000;
moduleToObjectTimer.clear();
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index f9269e0934271..a638f6e5eae7a 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -422,7 +422,9 @@ std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
// Translate the Module to ISA.
FailureOr<std::string> serializedISA =
- mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine);
+ mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine, [&]() {
+ return getOperation().emitError();
+ });
if (failed(serializedISA)) {
getOperation().emitError() << "failed translating the module to ISA";
return std::nullopt;
diff --git a/mlir/lib/Target/LLVM/XeVM/Target.cpp b/mlir/lib/Target/LLVM/XeVM/Target.cpp
index d78475cbd0c22..d1bd4d78d28c0 100644
--- a/mlir/lib/Target/LLVM/XeVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/XeVM/Target.cpp
@@ -302,7 +302,9 @@ SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
if (targetOptions.getCompilationTarget() ==
gpu::CompilationTarget::Assembly) {
FailureOr<std::string> serializedISA =
- mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine);
+ mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine, [&]() {
+ return getGPUModuleOp().emitError();
+ });
if (failed(serializedISA)) {
getGPUModuleOp().emitError() << "Failed translating the module to ISA."
<< triple << ", can't compile with LLVM\n";
>From 4add3dfa2a651317974ab160f1d31db14f56712a Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Sun, 14 Dec 2025 12:52:07 +0100
Subject: [PATCH 5/7] doc fix
---
mlir/include/mlir/Target/LLVM/ModuleToObject.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index 5be399065ef74..8c49e7712d6d3 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -24,7 +24,7 @@ class TargetMachine;
namespace mlir {
namespace LLVM {
-/// Utility function for translating to ISA.
+/// Translate LLVM module to textual ISA.
FailureOr<std::string>
translateModuleToISA(llvm::Module &llvmModule,
llvm::TargetMachine &targetMachine,
>From 47c290e210cb7831cb61c619e6a31034859ee5ee Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Sun, 14 Dec 2025 13:38:28 +0100
Subject: [PATCH 6/7] move translateModuleToISA
---
mlir/include/mlir/Target/LLVM/ModuleToObject.h | 13 +++++++------
mlir/lib/Target/LLVM/ModuleToObject.cpp | 7 +++----
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, 16 insertions(+), 19 deletions(-)
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index 8c49e7712d6d3..63b502598edaa 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -24,12 +24,6 @@ class TargetMachine;
namespace mlir {
namespace LLVM {
-/// Translate LLVM module to textual ISA.
-FailureOr<std::string>
-translateModuleToISA(llvm::Module &llvmModule,
- llvm::TargetMachine &targetMachine,
- function_ref<InFlightDiagnostic()> emitError);
-
class ModuleTranslation;
/// Utility base class for transforming operations into binary objects, by
/// default it returns the serialized LLVM bitcode for the module. The
@@ -51,6 +45,13 @@ class ModuleToObject {
/// Runs the serialization pipeline, returning `std::nullopt` on error.
virtual std::optional<SmallVector<char, 0>> run();
+ /// Translate LLVM module to textual ISA.
+ /// TODO: switch to SmallString
+ static FailureOr<std::string>
+ translateModuleToISA(llvm::Module &llvmModule,
+ llvm::TargetMachine &targetMachine,
+ function_ref<InFlightDiagnostic()> emitError);
+
protected:
// Hooks to be implemented by derived classes.
diff --git a/mlir/lib/Target/LLVM/ModuleToObject.cpp b/mlir/lib/Target/LLVM/ModuleToObject.cpp
index eb99e69b626b6..e20fd544ed8f7 100644
--- a/mlir/lib/Target/LLVM/ModuleToObject.cpp
+++ b/mlir/lib/Target/LLVM/ModuleToObject.cpp
@@ -205,10 +205,9 @@ LogicalResult ModuleToObject::optimizeModule(llvm::Module &module,
return success();
}
-FailureOr<std::string>
-mlir::LLVM::translateModuleToISA(llvm::Module &llvmModule,
- llvm::TargetMachine &targetMachine,
- function_ref<InFlightDiagnostic()> emitError) {
+FailureOr<std::string> ModuleToObject::translateModuleToISA(
+ llvm::Module &llvmModule, llvm::TargetMachine &targetMachine,
+ function_ref<InFlightDiagnostic()> emitError) {
std::string targetISA;
llvm::raw_string_ostream stream(targetISA);
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index c22f8a23d2b2b..cae5adfc086df 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -698,9 +698,8 @@ NVPTXSerializer::moduleToObject(llvm::Module &llvmModule) {
}
moduleToObjectTimer.startTimer();
FailureOr<std::string> serializedISA =
- mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine, [&]() {
- return getOperation().emitError();
- });
+ translateModuleToISA(llvmModule, **targetMachine,
+ [&]() { return getOperation().emitError(); });
moduleToObjectTimer.stopTimer();
llvmToISATimeInMs = moduleToObjectTimer.getTotalTime().getWallTime() * 1000;
moduleToObjectTimer.clear();
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index a638f6e5eae7a..a1970dac8785a 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -422,9 +422,8 @@ std::optional<SmallVector<char, 0>> SerializeGPUModuleBase::moduleToObjectImpl(
// Translate the Module to ISA.
FailureOr<std::string> serializedISA =
- mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine, [&]() {
- return getOperation().emitError();
- });
+ translateModuleToISA(llvmModule, **targetMachine,
+ [&]() { return getOperation().emitError(); });
if (failed(serializedISA)) {
getOperation().emitError() << "failed translating the module to ISA";
return std::nullopt;
diff --git a/mlir/lib/Target/LLVM/XeVM/Target.cpp b/mlir/lib/Target/LLVM/XeVM/Target.cpp
index d1bd4d78d28c0..24b518e6eed4c 100644
--- a/mlir/lib/Target/LLVM/XeVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/XeVM/Target.cpp
@@ -302,9 +302,8 @@ SPIRVSerializer::moduleToObject(llvm::Module &llvmModule) {
if (targetOptions.getCompilationTarget() ==
gpu::CompilationTarget::Assembly) {
FailureOr<std::string> serializedISA =
- mlir::LLVM::translateModuleToISA(llvmModule, **targetMachine, [&]() {
- return getGPUModuleOp().emitError();
- });
+ translateModuleToISA(llvmModule, **targetMachine,
+ [&]() { return getGPUModuleOp().emitError(); });
if (failed(serializedISA)) {
getGPUModuleOp().emitError() << "Failed translating the module to ISA."
<< triple << ", can't compile with LLVM\n";
>From f1ec22214b5432fc3b2e1be8679cc1c2b794bd6a Mon Sep 17 00:00:00 2001
From: Ivan Butygin <ivan.butygin at gmail.com>
Date: Sun, 14 Dec 2025 13:40:39 +0100
Subject: [PATCH 7/7] stray newline
---
mlir/include/mlir/Target/LLVM/ModuleToObject.h | 1 -
1 file changed, 1 deletion(-)
diff --git a/mlir/include/mlir/Target/LLVM/ModuleToObject.h b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
index 63b502598edaa..ea0b4b4a6ab7f 100644
--- a/mlir/include/mlir/Target/LLVM/ModuleToObject.h
+++ b/mlir/include/mlir/Target/LLVM/ModuleToObject.h
@@ -23,7 +23,6 @@ class TargetMachine;
namespace mlir {
namespace LLVM {
-
class ModuleTranslation;
/// Utility base class for transforming operations into binary objects, by
/// default it returns the serialized LLVM bitcode for the module. The
More information about the Mlir-commits
mailing list