[Mlir-commits] [mlir] [WIP]: [mlir][gpu] Adding ELF section option to the gpu-module-to-binary pass (PR #119734)
Renaud Kauffmann
llvmlistbot at llvm.org
Thu Dec 12 10:18:31 PST 2024
https://github.com/Renaud-K created https://github.com/llvm/llvm-project/pull/119734
#119440
>From 6d3cb91fa6d24665a931b17225146072a18c2579 Mon Sep 17 00:00:00 2001
From: Renaud-K <rkauffmann at nvidia.com>
Date: Tue, 10 Dec 2024 11:48:45 -0800
Subject: [PATCH 1/3] Adding the ELF section to the NVVMTargetAttr, to
propagate to the gpu.binary Op
---
mlir/include/mlir/Dialect/GPU/Transforms/Passes.td | 3 +++
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 6 ++++--
mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp | 4 +++-
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 4 ++--
mlir/lib/Target/LLVM/NVVM/Target.cpp | 12 ++++++++++--
mlir/test/Dialect/LLVMIR/attach-targets.mlir | 3 +++
6 files changed, 25 insertions(+), 7 deletions(-)
diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
index 4a9ddafdd177d2..784721d19d0ccd 100644
--- a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
@@ -139,6 +139,9 @@ def GpuNVVMAttachTarget: Pass<"nvvm-attach-target", ""> {
Option<"ftzFlag", "ftz", "bool",
/*default=*/"false",
"Enable flush to zero for denormals.">,
+ Option<"elfSection", "section", "std::string",
+ /*default=*/"\"\"",
+ "ELF section where the module needs to be created.">,
ListOption<"linkLibs", "l", "std::string",
"Extra bitcode libraries paths to link to.">,
];
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 14880a1a66ba57..747a711eefcd96 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -2367,20 +2367,22 @@ def NVVM_TargettAttr : NVVM_Attr<"NVVMTarget", "target"> {
StringRefParameter<"Target triple.", "\"nvptx64-nvidia-cuda\"">:$triple,
StringRefParameter<"Target chip.", "\"sm_50\"">:$chip,
StringRefParameter<"Target chip features.", "\"+ptx60\"">:$features,
+ OptionalParameter<"StringAttr", "ELF section.">:$section,
OptionalParameter<"DictionaryAttr", "Target specific flags.">:$flags,
OptionalParameter<"ArrayAttr", "Files to link to the LLVM module.">:$link
);
let assemblyFormat = [{
- (`<` struct($O, $triple, $chip, $features, $flags, $link)^ `>`)?
+ (`<` struct($O, $triple, $chip, $features, $section, $flags, $link)^ `>`)?
}];
let builders = [
AttrBuilder<(ins CArg<"int", "2">:$optLevel,
CArg<"StringRef", "\"nvptx64-nvidia-cuda\"">:$triple,
CArg<"StringRef", "\"sm_50\"">:$chip,
CArg<"StringRef", "\"+ptx60\"">:$features,
+ CArg<"StringAttr", "nullptr">:$section,
CArg<"DictionaryAttr", "nullptr">:$targetFlags,
CArg<"ArrayAttr", "nullptr">:$linkFiles), [{
- return Base::get($_ctxt, optLevel, triple, chip, features, targetFlags, linkFiles);
+ return Base::get($_ctxt, optLevel, triple, chip, features, section, targetFlags, linkFiles);
}]>
];
let skipDefaultBuilders = 1;
diff --git a/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp b/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
index dd705cd3383121..5628361c6a843c 100644
--- a/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
@@ -63,7 +63,9 @@ void NVVMAttachTarget::runOnOperation() {
ArrayRef<std::string> libs(linkLibs);
SmallVector<StringRef> filesToLink(libs);
auto target = builder.getAttr<NVVMTargetAttr>(
- optLevel, triple, chip, features, getFlags(builder),
+ optLevel, triple, chip, features,
+ elfSection.empty() ? nullptr : builder.getStringAttr(elfSection),
+ getFlags(builder),
filesToLink.empty() ? nullptr : builder.getStrArrayAttr(filesToLink));
llvm::Regex matcher(moduleMatcher);
for (Region ®ion : getOperation()->getRegions())
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 8b09c0f386d6b6..298e36e138c8e1 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -1249,8 +1249,8 @@ LogicalResult NVVMDialect::verifyRegionArgAttribute(Operation *op,
LogicalResult
NVVMTargetAttr::verify(function_ref<InFlightDiagnostic()> emitError,
int optLevel, StringRef triple, StringRef chip,
- StringRef features, DictionaryAttr flags,
- ArrayAttr files) {
+ StringRef features, StringAttr elfSection,
+ DictionaryAttr flags, ArrayAttr files) {
if (optLevel < 0 || optLevel > 3) {
emitError() << "The optimization level must be a number between 0 and 3.";
return failure();
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 3c92359915ded4..152d2dd6326666 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -664,9 +664,17 @@ NVVMTargetAttrImpl::createObject(Attribute attribute, Operation *module,
gpu::CompilationTarget format = options.getCompilationTarget();
DictionaryAttr objectProps;
Builder builder(attribute.getContext());
+ SmallVector<NamedAttribute, 2> properties;
if (format == gpu::CompilationTarget::Assembly)
- objectProps = builder.getDictionaryAttr(
- {builder.getNamedAttr("O", builder.getI32IntegerAttr(target.getO()))});
+ properties.push_back(
+ builder.getNamedAttr("O", builder.getI32IntegerAttr(target.getO())));
+
+ if (StringAttr section = target.getSection())
+ properties.push_back(builder.getNamedAttr("section", section));
+
+ if (!properties.empty())
+ objectProps = builder.getDictionaryAttr(properties);
+
return builder.getAttr<gpu::ObjectAttr>(
attribute, format,
builder.getStringAttr(StringRef(object.data(), object.size())),
diff --git a/mlir/test/Dialect/LLVMIR/attach-targets.mlir b/mlir/test/Dialect/LLVMIR/attach-targets.mlir
index 83733db400798e..832e0443e720f7 100644
--- a/mlir/test/Dialect/LLVMIR/attach-targets.mlir
+++ b/mlir/test/Dialect/LLVMIR/attach-targets.mlir
@@ -1,13 +1,16 @@
// RUN: mlir-opt %s --nvvm-attach-target='module=nvvm.* O=3 chip=sm_90' --rocdl-attach-target='module=rocdl.* O=3 chip=gfx90a' | FileCheck %s
// RUN: mlir-opt %s --nvvm-attach-target='module=options.* O=1 chip=sm_70 fast=true ftz=true' --rocdl-attach-target='module=options.* l=file1.bc,file2.bc wave64=false finite-only=true' | FileCheck %s --check-prefix=CHECK_OPTS
+// RUN: mlir-opt %s --nvvm-attach-target='module=nvvm.* section=__fatbin' | FileCheck %s --check-prefix=CHECK_SECTION
module attributes {gpu.container_module} {
// Verify the target is appended.
// CHECK: @nvvm_module_1 [#nvvm.target<O = 3, chip = "sm_90">] {
+// CHECK_SECTION: @nvvm_module_1 [#nvvm.target<section = "__fatbin">]
gpu.module @nvvm_module_1 {
}
// Verify the target is appended.
// CHECK: @nvvm_module_2 [#nvvm.target<chip = "sm_60">, #nvvm.target<O = 3, chip = "sm_90">] {
+// CHECK_SECTION: gpu.module @nvvm_module_2 [#nvvm.target<chip = "sm_60">, #nvvm.target<section = "__fatbin">]
gpu.module @nvvm_module_2 [#nvvm.target<chip = "sm_60">] {
}
// Verify the target is not added multiple times.
>From 8b1c249b537cc08076d23b6e05515e65d0b071e3 Mon Sep 17 00:00:00 2001
From: Renaud-K <rkauffmann at nvidia.com>
Date: Tue, 10 Dec 2024 15:26:21 -0800
Subject: [PATCH 2/3] Adding ELF section option to the gpu-module-to-binary
pass
---
.../Dialect/GPU/IR/CompilationInterfaces.h | 9 ++++++++-
.../mlir/Dialect/GPU/Transforms/Passes.td | 19 +++++++++----------
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp | 18 ++++++++++++------
.../Dialect/GPU/Transforms/ModuleToBinary.cpp | 4 ++--
mlir/lib/Target/LLVM/NVVM/Target.cpp | 5 +++--
.../Dialect/GPU/module-to-binary-nvvm.mlir | 3 +++
.../Target/LLVM/SerializeNVVMTarget.cpp | 12 ++++++------
.../Target/LLVM/SerializeROCDLTarget.cpp | 12 ++++++------
.../Target/LLVM/SerializeToLLVMBitcode.cpp | 12 ++++++------
9 files changed, 55 insertions(+), 39 deletions(-)
diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
index d4b16a1de8eddc..8cf8a7e55517d4 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
+++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
@@ -51,7 +51,7 @@ class TargetOptions {
/// `Fatbin`.
TargetOptions(
StringRef toolkitPath = {}, ArrayRef<std::string> linkFiles = {},
- StringRef cmdOptions = {},
+ StringRef cmdOptions = {}, StringRef elfSection = {},
CompilationTarget compilationTarget = getDefaultCompilationTarget(),
function_ref<SymbolTable *()> getSymbolTableCallback = {},
function_ref<void(llvm::Module &)> initialLlvmIRCallback = {},
@@ -71,6 +71,9 @@ class TargetOptions {
/// Returns the command line options.
StringRef getCmdOptions() const;
+ /// Returns the ELF section.
+ StringRef getELFSection() const;
+
/// Returns a tokenization of the command line options.
std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
tokenizeCmdOptions() const;
@@ -110,6 +113,7 @@ class TargetOptions {
TargetOptions(
TypeID typeID, StringRef toolkitPath = {},
ArrayRef<std::string> linkFiles = {}, StringRef cmdOptions = {},
+ StringRef elfSection = {},
CompilationTarget compilationTarget = getDefaultCompilationTarget(),
function_ref<SymbolTable *()> getSymbolTableCallback = {},
function_ref<void(llvm::Module &)> initialLlvmIRCallback = {},
@@ -127,6 +131,9 @@ class TargetOptions {
/// process.
std::string cmdOptions;
+ /// ELF Section where the binary needs to be located
+ std::string elfSection;
+
/// Compilation process target format.
CompilationTarget compilationTarget;
diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
index 784721d19d0ccd..cadfd866cb259d 100644
--- a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
@@ -87,16 +87,15 @@ def GpuModuleToBinaryPass
3. `binary`, `bin`: produces binaries.
4. `fatbinary`, `fatbin`: produces fatbinaries.
}];
- let options = [
- Option<"toolkitPath", "toolkit", "std::string", [{""}],
- "Toolkit path.">,
- ListOption<"linkFiles", "l", "std::string",
- "Extra files to link to.">,
- Option<"cmdOptions", "opts", "std::string", [{""}],
- "Command line options to pass to the tools.">,
- Option<"compilationTarget", "format", "std::string", [{"fatbin"}],
- "The target representation of the compilation process.">
- ];
+ let options =
+ [Option<"toolkitPath", "toolkit", "std::string", [{""}], "Toolkit path.">,
+ ListOption<"linkFiles", "l", "std::string", "Extra files to link to.">,
+ Option<"cmdOptions", "opts", "std::string", [{""}],
+ "Command line options to pass to the tools.">,
+ Option<"compilationTarget", "format", "std::string", [{"fatbin"}],
+ "The target representation of the compilation process.">,
+ Option<"elfSection", "section", "std::string", [{""}],
+ "ELF section where binary is to be located.">];
}
def GpuNVVMAttachTarget: Pass<"nvvm-attach-target", ""> {
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index ee00fbeb28b61d..1fad251b2f79e0 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -2484,27 +2484,31 @@ KernelMetadataAttr KernelTableAttr::lookup(StringAttr key) const {
TargetOptions::TargetOptions(
StringRef toolkitPath, ArrayRef<std::string> linkFiles,
- StringRef cmdOptions, CompilationTarget compilationTarget,
+ StringRef cmdOptions, StringRef elfSection,
+ CompilationTarget compilationTarget,
function_ref<SymbolTable *()> getSymbolTableCallback,
function_ref<void(llvm::Module &)> initialLlvmIRCallback,
function_ref<void(llvm::Module &)> linkedLlvmIRCallback,
function_ref<void(llvm::Module &)> optimizedLlvmIRCallback,
function_ref<void(StringRef)> isaCallback)
: TargetOptions(TypeID::get<TargetOptions>(), toolkitPath, linkFiles,
- cmdOptions, compilationTarget, getSymbolTableCallback,
- initialLlvmIRCallback, linkedLlvmIRCallback,
- optimizedLlvmIRCallback, isaCallback) {}
+ cmdOptions, elfSection, compilationTarget,
+ getSymbolTableCallback, initialLlvmIRCallback,
+ linkedLlvmIRCallback, optimizedLlvmIRCallback,
+ isaCallback) {}
TargetOptions::TargetOptions(
TypeID typeID, StringRef toolkitPath, ArrayRef<std::string> linkFiles,
- StringRef cmdOptions, CompilationTarget compilationTarget,
+ StringRef cmdOptions, StringRef elfSection,
+ CompilationTarget compilationTarget,
function_ref<SymbolTable *()> getSymbolTableCallback,
function_ref<void(llvm::Module &)> initialLlvmIRCallback,
function_ref<void(llvm::Module &)> linkedLlvmIRCallback,
function_ref<void(llvm::Module &)> optimizedLlvmIRCallback,
function_ref<void(StringRef)> isaCallback)
: toolkitPath(toolkitPath.str()), linkFiles(linkFiles),
- cmdOptions(cmdOptions.str()), compilationTarget(compilationTarget),
+ cmdOptions(cmdOptions.str()), elfSection(elfSection.str()),
+ compilationTarget(compilationTarget),
getSymbolTableCallback(getSymbolTableCallback),
initialLlvmIRCallback(initialLlvmIRCallback),
linkedLlvmIRCallback(linkedLlvmIRCallback),
@@ -2519,6 +2523,8 @@ ArrayRef<std::string> TargetOptions::getLinkFiles() const { return linkFiles; }
StringRef TargetOptions::getCmdOptions() const { return cmdOptions; }
+StringRef TargetOptions::getELFSection() const { return elfSection; }
+
SymbolTable *TargetOptions::getSymbolTable() const {
return getSymbolTableCallback ? getSymbolTableCallback() : nullptr;
}
diff --git a/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp b/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
index 86a3b4780e88ce..295ece4782fdbf 100644
--- a/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
@@ -69,8 +69,8 @@ void GpuModuleToBinaryPass::runOnOperation() {
return &parentTable.value();
};
- TargetOptions targetOptions(toolkitPath, linkFiles, cmdOptions, *targetFormat,
- lazyTableBuilder);
+ TargetOptions targetOptions(toolkitPath, linkFiles, cmdOptions, elfSection,
+ *targetFormat, lazyTableBuilder);
if (failed(transformGpuModulesToBinaries(
getOperation(), OffloadingLLVMTranslationAttrInterface(nullptr),
targetOptions)))
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 152d2dd6326666..e39c63d2b422bb 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -669,8 +669,9 @@ NVVMTargetAttrImpl::createObject(Attribute attribute, Operation *module,
properties.push_back(
builder.getNamedAttr("O", builder.getI32IntegerAttr(target.getO())));
- if (StringAttr section = target.getSection())
- properties.push_back(builder.getNamedAttr("section", section));
+ if (auto section = options.getELFSection(); !section.empty())
+ properties.push_back(
+ builder.getNamedAttr("section", builder.getStringAttr(section)));
if (!properties.empty())
objectProps = builder.getDictionaryAttr(properties);
diff --git a/mlir/test/Dialect/GPU/module-to-binary-nvvm.mlir b/mlir/test/Dialect/GPU/module-to-binary-nvvm.mlir
index c286c8bc9042ff..8b93e7ef983a6c 100644
--- a/mlir/test/Dialect/GPU/module-to-binary-nvvm.mlir
+++ b/mlir/test/Dialect/GPU/module-to-binary-nvvm.mlir
@@ -1,10 +1,12 @@
// REQUIRES: host-supports-nvptx
// RUN: mlir-opt %s --gpu-module-to-binary="format=llvm" | FileCheck %s
// RUN: mlir-opt %s --gpu-module-to-binary="format=isa" | FileCheck %s -check-prefix=CHECK-ISA
+// RUN: mlir-opt %s --gpu-module-to-binary="section=__fatbin" | FileCheck %s -check-prefix=CHECK-SECTION
module attributes {gpu.container_module} {
// CHECK-LABEL:gpu.binary @kernel_module1
// CHECK:[#gpu.object<#nvvm.target<chip = "sm_70">, offload = "{{.*}}">]
+ // CHECK-SECTION: #gpu.object<#nvvm.target<chip = "sm_70">, properties = {section = "__fatbin"}
gpu.module @kernel_module1 [#nvvm.target<chip = "sm_70">] {
llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
%arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
@@ -25,6 +27,7 @@ module attributes {gpu.container_module} {
// CHECK-LABEL:gpu.binary @kernel_module3 <#gpu.select_object<1 : i64>>
// CHECK:[#gpu.object<#nvvm.target<chip = "sm_70">, offload = "{{.*}}">, #gpu.object<#nvvm.target<chip = "sm_80">, offload = "{{.*}}">]
+ // CHECK-SECTION: [#gpu.object<#nvvm.target<chip = "sm_70">, properties = {section = "__fatbin"},{{.*}} #gpu.object<#nvvm.target<chip = "sm_80">, properties = {section = "__fatbin"}
gpu.module @kernel_module3 <#gpu.select_object<1>> [
#nvvm.target<chip = "sm_70">,
#nvvm.target<chip = "sm_80">] {
diff --git a/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp b/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
index eee9bd5f234751..a92ad18c956821 100644
--- a/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
+++ b/mlir/unittests/Target/LLVM/SerializeNVVMTarget.cpp
@@ -81,7 +81,7 @@ TEST_F(MLIRTargetLLVMNVVM, SKIP_WITHOUT_NVPTX(SerializeNVVMMToLLVM)) {
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Offload);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Offload);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
@@ -117,7 +117,7 @@ TEST_F(MLIRTargetLLVMNVVM, SKIP_WITHOUT_NVPTX(SerializeNVVMToPTX)) {
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Assembly);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Assembly);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
@@ -147,7 +147,7 @@ TEST_F(MLIRTargetLLVMNVVM, SKIP_WITHOUT_NVPTX(SerializeNVVMToBinary)) {
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Binary);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Binary);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
@@ -194,9 +194,9 @@ TEST_F(MLIRTargetLLVMNVVM,
isaResult = isa.str();
};
- gpu::TargetOptions options({}, {}, {}, gpu::CompilationTarget::Assembly, {},
- initialCallback, linkedCallback, optimizedCallback,
- isaCallback);
+ gpu::TargetOptions options({}, {}, {}, {}, gpu::CompilationTarget::Assembly,
+ {}, initialCallback, linkedCallback,
+ optimizedCallback, isaCallback);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
diff --git a/mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp b/mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp
index b02f34c812b3f4..d5e72a17131748 100644
--- a/mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp
+++ b/mlir/unittests/Target/LLVM/SerializeROCDLTarget.cpp
@@ -83,7 +83,7 @@ TEST_F(MLIRTargetLLVMROCDL, SKIP_WITHOUT_AMDGPU(SerializeROCDLToLLVM)) {
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Offload);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Offload);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
@@ -119,7 +119,7 @@ TEST_F(MLIRTargetLLVMROCDL,
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Assembly);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Assembly);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
@@ -145,7 +145,7 @@ TEST_F(MLIRTargetLLVMROCDL,
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Assembly);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Assembly);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
@@ -169,7 +169,7 @@ TEST_F(MLIRTargetLLVMROCDL, SKIP_WITHOUT_AMDGPU(SerializeROCDLToPTX)) {
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Assembly);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Assembly);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
@@ -199,7 +199,7 @@ TEST_F(MLIRTargetLLVMROCDL, SKIP_WITHOUT_AMDGPU(SerializeROCDLToBinary)) {
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Binary);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Binary);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
@@ -243,7 +243,7 @@ TEST_F(MLIRTargetLLVMROCDL, SKIP_WITHOUT_AMDGPU(GetELFMetadata)) {
// Serialize the module.
auto serializer = dyn_cast<gpu::TargetAttrInterface>(target);
ASSERT_TRUE(!!serializer);
- gpu::TargetOptions options("", {}, "", gpu::CompilationTarget::Binary);
+ gpu::TargetOptions options("", {}, "", "", gpu::CompilationTarget::Binary);
for (auto gpuModule : (*module).getBody()->getOps<gpu::GPUModuleOp>()) {
std::optional<SmallVector<char, 0>> object =
serializer.serializeToObject(gpuModule, options);
diff --git a/mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp b/mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp
index 63d1dbd2519bea..314f25833f38ca 100644
--- a/mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp
+++ b/mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp
@@ -174,8 +174,8 @@ TEST_F(MLIRTargetLLVM, SKIP_WITHOUT_NATIVE(CallbackInvokedWithInitialLLVMIR)) {
};
gpu::TargetOptions opts(
- {}, {}, {}, mlir::gpu::TargetOptions::getDefaultCompilationTarget(), {},
- initialCallback);
+ {}, {}, {}, {}, mlir::gpu::TargetOptions::getDefaultCompilationTarget(),
+ {}, initialCallback);
std::optional<SmallVector<char, 0>> serializedBinary =
targetAttr.serializeToObject(*module, opts);
@@ -202,8 +202,8 @@ TEST_F(MLIRTargetLLVM, SKIP_WITHOUT_NATIVE(CallbackInvokedWithLinkedLLVMIR)) {
};
gpu::TargetOptions opts(
- {}, {}, {}, mlir::gpu::TargetOptions::getDefaultCompilationTarget(), {},
- {}, linkedCallback);
+ {}, {}, {}, {}, mlir::gpu::TargetOptions::getDefaultCompilationTarget(),
+ {}, {}, linkedCallback);
std::optional<SmallVector<char, 0>> serializedBinary =
targetAttr.serializeToObject(*module, opts);
@@ -231,8 +231,8 @@ TEST_F(MLIRTargetLLVM,
};
gpu::TargetOptions opts(
- {}, {}, {}, mlir::gpu::TargetOptions::getDefaultCompilationTarget(), {},
- {}, {}, optimizedCallback);
+ {}, {}, {}, {}, mlir::gpu::TargetOptions::getDefaultCompilationTarget(),
+ {}, {}, {}, optimizedCallback);
std::optional<SmallVector<char, 0>> serializedBinary =
targetAttr.serializeToObject(*module, opts);
>From be020390fea8319f49d58c37440d859b66061f9e Mon Sep 17 00:00:00 2001
From: Renaud-K <rkauffmann at nvidia.com>
Date: Tue, 10 Dec 2024 15:39:48 -0800
Subject: [PATCH 3/3] Reverting initial commit
---
mlir/include/mlir/Dialect/GPU/Transforms/Passes.td | 3 ---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 6 ++----
mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp | 4 +---
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 4 ++--
mlir/test/Dialect/LLVMIR/attach-targets.mlir | 3 ---
5 files changed, 5 insertions(+), 15 deletions(-)
diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
index cadfd866cb259d..e9e5dade84c988 100644
--- a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
@@ -138,9 +138,6 @@ def GpuNVVMAttachTarget: Pass<"nvvm-attach-target", ""> {
Option<"ftzFlag", "ftz", "bool",
/*default=*/"false",
"Enable flush to zero for denormals.">,
- Option<"elfSection", "section", "std::string",
- /*default=*/"\"\"",
- "ELF section where the module needs to be created.">,
ListOption<"linkLibs", "l", "std::string",
"Extra bitcode libraries paths to link to.">,
];
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 747a711eefcd96..14880a1a66ba57 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -2367,22 +2367,20 @@ def NVVM_TargettAttr : NVVM_Attr<"NVVMTarget", "target"> {
StringRefParameter<"Target triple.", "\"nvptx64-nvidia-cuda\"">:$triple,
StringRefParameter<"Target chip.", "\"sm_50\"">:$chip,
StringRefParameter<"Target chip features.", "\"+ptx60\"">:$features,
- OptionalParameter<"StringAttr", "ELF section.">:$section,
OptionalParameter<"DictionaryAttr", "Target specific flags.">:$flags,
OptionalParameter<"ArrayAttr", "Files to link to the LLVM module.">:$link
);
let assemblyFormat = [{
- (`<` struct($O, $triple, $chip, $features, $section, $flags, $link)^ `>`)?
+ (`<` struct($O, $triple, $chip, $features, $flags, $link)^ `>`)?
}];
let builders = [
AttrBuilder<(ins CArg<"int", "2">:$optLevel,
CArg<"StringRef", "\"nvptx64-nvidia-cuda\"">:$triple,
CArg<"StringRef", "\"sm_50\"">:$chip,
CArg<"StringRef", "\"+ptx60\"">:$features,
- CArg<"StringAttr", "nullptr">:$section,
CArg<"DictionaryAttr", "nullptr">:$targetFlags,
CArg<"ArrayAttr", "nullptr">:$linkFiles), [{
- return Base::get($_ctxt, optLevel, triple, chip, features, section, targetFlags, linkFiles);
+ return Base::get($_ctxt, optLevel, triple, chip, features, targetFlags, linkFiles);
}]>
];
let skipDefaultBuilders = 1;
diff --git a/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp b/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
index 5628361c6a843c..dd705cd3383121 100644
--- a/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
@@ -63,9 +63,7 @@ void NVVMAttachTarget::runOnOperation() {
ArrayRef<std::string> libs(linkLibs);
SmallVector<StringRef> filesToLink(libs);
auto target = builder.getAttr<NVVMTargetAttr>(
- optLevel, triple, chip, features,
- elfSection.empty() ? nullptr : builder.getStringAttr(elfSection),
- getFlags(builder),
+ optLevel, triple, chip, features, getFlags(builder),
filesToLink.empty() ? nullptr : builder.getStrArrayAttr(filesToLink));
llvm::Regex matcher(moduleMatcher);
for (Region ®ion : getOperation()->getRegions())
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 298e36e138c8e1..8b09c0f386d6b6 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -1249,8 +1249,8 @@ LogicalResult NVVMDialect::verifyRegionArgAttribute(Operation *op,
LogicalResult
NVVMTargetAttr::verify(function_ref<InFlightDiagnostic()> emitError,
int optLevel, StringRef triple, StringRef chip,
- StringRef features, StringAttr elfSection,
- DictionaryAttr flags, ArrayAttr files) {
+ StringRef features, DictionaryAttr flags,
+ ArrayAttr files) {
if (optLevel < 0 || optLevel > 3) {
emitError() << "The optimization level must be a number between 0 and 3.";
return failure();
diff --git a/mlir/test/Dialect/LLVMIR/attach-targets.mlir b/mlir/test/Dialect/LLVMIR/attach-targets.mlir
index 832e0443e720f7..83733db400798e 100644
--- a/mlir/test/Dialect/LLVMIR/attach-targets.mlir
+++ b/mlir/test/Dialect/LLVMIR/attach-targets.mlir
@@ -1,16 +1,13 @@
// RUN: mlir-opt %s --nvvm-attach-target='module=nvvm.* O=3 chip=sm_90' --rocdl-attach-target='module=rocdl.* O=3 chip=gfx90a' | FileCheck %s
// RUN: mlir-opt %s --nvvm-attach-target='module=options.* O=1 chip=sm_70 fast=true ftz=true' --rocdl-attach-target='module=options.* l=file1.bc,file2.bc wave64=false finite-only=true' | FileCheck %s --check-prefix=CHECK_OPTS
-// RUN: mlir-opt %s --nvvm-attach-target='module=nvvm.* section=__fatbin' | FileCheck %s --check-prefix=CHECK_SECTION
module attributes {gpu.container_module} {
// Verify the target is appended.
// CHECK: @nvvm_module_1 [#nvvm.target<O = 3, chip = "sm_90">] {
-// CHECK_SECTION: @nvvm_module_1 [#nvvm.target<section = "__fatbin">]
gpu.module @nvvm_module_1 {
}
// Verify the target is appended.
// CHECK: @nvvm_module_2 [#nvvm.target<chip = "sm_60">, #nvvm.target<O = 3, chip = "sm_90">] {
-// CHECK_SECTION: gpu.module @nvvm_module_2 [#nvvm.target<chip = "sm_60">, #nvvm.target<section = "__fatbin">]
gpu.module @nvvm_module_2 [#nvvm.target<chip = "sm_60">] {
}
// Verify the target is not added multiple times.
More information about the Mlir-commits
mailing list