[mlir][nvvm_target] Add nvvm options into NVVMAttachTarget pass #126478
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Member
|
@llvm/pr-subscribers-mlir Author: Zichen Lu (MikaOvO) ChangesAdd Then user can pass it to something (such as libnvvm.so) when lowering the GpuModule to binary. Full diff: https://github.com/llvm/llvm-project/pull/126478.diff 4 Files Affected:
diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
index c950ef220f692f6..9a890ae24d8fc37 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
+++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
@@ -108,6 +108,10 @@ class TargetOptions {
/// Returns the default compilation target: `CompilationTarget::Fatbin`.
static CompilationTarget getDefaultCompilationTarget();
+ /// Returns a tokenization of the command line options.
+ static std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+ tokenizeCmdOptions(const std::string &cmdOptions);
+
protected:
/// Derived classes must use this constructor to initialize `typeID` to the
/// appropiate value: ie. `TargetOptions(TypeID::get<DerivedClass>())`.
diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
index e055164a1c384ee..86459c506cf1055 100644
--- a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
@@ -141,6 +141,9 @@ def GpuNVVMAttachTarget: Pass<"nvvm-attach-target", ""> {
Option<"ftzFlag", "ftz", "bool",
/*default=*/"false",
"Enable flush to zero for denormals.">,
+ Option<"nvvmOptions", "nvvm-options", "std::string",
+ /*default=*/ [{""}],
+ "Options passed to libnvvm">,
ListOption<"linkLibs", "l", "std::string",
"Extra bitcode libraries paths to link to.">,
];
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index d06f10d3137a183..1bdeb3e356f4be9 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -2564,7 +2564,7 @@ CompilationTarget TargetOptions::getDefaultCompilationTarget() {
}
std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
-TargetOptions::tokenizeCmdOptions() const {
+TargetOptions::tokenizeCmdOptions(const std::string &cmdOptions) {
std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>> options;
llvm::StringSaver stringSaver(options.first);
StringRef opts = cmdOptions;
@@ -2586,6 +2586,11 @@ TargetOptions::tokenizeCmdOptions() const {
return options;
}
+std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+TargetOptions::tokenizeCmdOptions() const {
+ return tokenizeCmdOptions(cmdOptions);
+}
+
MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions)
#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc"
diff --git a/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp b/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
index dd705cd33831217..663bcbf7a848cea 100644
--- a/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
@@ -45,7 +45,7 @@ struct NVVMAttachTarget
DictionaryAttr NVVMAttachTarget::getFlags(OpBuilder &builder) const {
UnitAttr unitAttr = builder.getUnitAttr();
- SmallVector<NamedAttribute, 2> flags;
+ SmallVector<NamedAttribute, 3> flags;
auto addFlag = [&](StringRef flag) {
flags.push_back(builder.getNamedAttr(flag, unitAttr));
};
@@ -53,6 +53,19 @@ DictionaryAttr NVVMAttachTarget::getFlags(OpBuilder &builder) const {
addFlag("fast");
if (ftzFlag)
addFlag("ftz");
+ if (!nvvmOptions.empty()) {
+ auto options = gpu::TargetOptions::tokenizeCmdOptions(nvvmOptions);
+ if (!options.second.empty()) {
+ llvm::SmallVector<mlir::Attribute> nvvmOptionAttrs;
+ for (const char *opt : options.second) {
+ nvvmOptionAttrs.emplace_back(
+ mlir::StringAttr::get(builder.getContext(), StringRef(opt)));
+ }
+ flags.push_back(builder.getNamedAttr(
+ "libNVVMOptions",
+ mlir::ArrayAttr::get(builder.getContext(), nvvmOptionAttrs)));
+ }
+ }
if (!flags.empty())
return builder.getDictionaryAttr(flags);
return nullptr;
|
Member
|
@llvm/pr-subscribers-mlir-gpu Author: Zichen Lu (MikaOvO) ChangesAdd Then user can pass it to something (such as libnvvm.so) when lowering the GpuModule to binary. Full diff: https://github.com/llvm/llvm-project/pull/126478.diff 4 Files Affected:
diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
index c950ef220f692f6..9a890ae24d8fc37 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
+++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationInterfaces.h
@@ -108,6 +108,10 @@ class TargetOptions {
/// Returns the default compilation target: `CompilationTarget::Fatbin`.
static CompilationTarget getDefaultCompilationTarget();
+ /// Returns a tokenization of the command line options.
+ static std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+ tokenizeCmdOptions(const std::string &cmdOptions);
+
protected:
/// Derived classes must use this constructor to initialize `typeID` to the
/// appropiate value: ie. `TargetOptions(TypeID::get<DerivedClass>())`.
diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
index e055164a1c384ee..86459c506cf1055 100644
--- a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
@@ -141,6 +141,9 @@ def GpuNVVMAttachTarget: Pass<"nvvm-attach-target", ""> {
Option<"ftzFlag", "ftz", "bool",
/*default=*/"false",
"Enable flush to zero for denormals.">,
+ Option<"nvvmOptions", "nvvm-options", "std::string",
+ /*default=*/ [{""}],
+ "Options passed to libnvvm">,
ListOption<"linkLibs", "l", "std::string",
"Extra bitcode libraries paths to link to.">,
];
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index d06f10d3137a183..1bdeb3e356f4be9 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -2564,7 +2564,7 @@ CompilationTarget TargetOptions::getDefaultCompilationTarget() {
}
std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
-TargetOptions::tokenizeCmdOptions() const {
+TargetOptions::tokenizeCmdOptions(const std::string &cmdOptions) {
std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>> options;
llvm::StringSaver stringSaver(options.first);
StringRef opts = cmdOptions;
@@ -2586,6 +2586,11 @@ TargetOptions::tokenizeCmdOptions() const {
return options;
}
+std::pair<llvm::BumpPtrAllocator, SmallVector<const char *>>
+TargetOptions::tokenizeCmdOptions() const {
+ return tokenizeCmdOptions(cmdOptions);
+}
+
MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::TargetOptions)
#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc"
diff --git a/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp b/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
index dd705cd33831217..663bcbf7a848cea 100644
--- a/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/NVVMAttachTarget.cpp
@@ -45,7 +45,7 @@ struct NVVMAttachTarget
DictionaryAttr NVVMAttachTarget::getFlags(OpBuilder &builder) const {
UnitAttr unitAttr = builder.getUnitAttr();
- SmallVector<NamedAttribute, 2> flags;
+ SmallVector<NamedAttribute, 3> flags;
auto addFlag = [&](StringRef flag) {
flags.push_back(builder.getNamedAttr(flag, unitAttr));
};
@@ -53,6 +53,19 @@ DictionaryAttr NVVMAttachTarget::getFlags(OpBuilder &builder) const {
addFlag("fast");
if (ftzFlag)
addFlag("ftz");
+ if (!nvvmOptions.empty()) {
+ auto options = gpu::TargetOptions::tokenizeCmdOptions(nvvmOptions);
+ if (!options.second.empty()) {
+ llvm::SmallVector<mlir::Attribute> nvvmOptionAttrs;
+ for (const char *opt : options.second) {
+ nvvmOptionAttrs.emplace_back(
+ mlir::StringAttr::get(builder.getContext(), StringRef(opt)));
+ }
+ flags.push_back(builder.getNamedAttr(
+ "libNVVMOptions",
+ mlir::ArrayAttr::get(builder.getContext(), nvvmOptionAttrs)));
+ }
+ }
if (!flags.empty())
return builder.getDictionaryAttr(flags);
return nullptr;
|
Member
|
MikaOvO
force-pushed
the
zilu/add_nvvm_options_in_nvvm_target
branch
2 times, most recently
from
0e0db64 to
8e3d841
Compare
MikaOvO
force-pushed
the
zilu/add_nvvm_options_in_nvvm_target
branch
from
8e3d841 to
f51b6a1
Compare
Contributor
Author
|
Hi, I have added a unittest. |
Member
|
Hi @MikaOvO why did we close this PR? I think it's useful to have it |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Add
nvvmOptionsinto NVVMAttachTarget pass.Then user can pass it to something (such as libnvvm.so) when lowering the GpuModule to binary.