[mlir][gpu] Pass GPU module to TargetAttrInterface::createObject.
#94910
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
|
@llvm/pr-subscribers-mlir-spirv @llvm/pr-subscribers-mlir-llvm Author: Fabian Mora (fabianmcg) ChangesThis patch adds an argument to Full diff: https://github.com/llvm/llvm-project/pull/94910.diff 5 Files Affected:
diff --git a/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td b/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
index 1607304803942..3ee4b215f7d64 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td
@@ -46,7 +46,8 @@ def GPUTargetAttrInterface : AttrInterface<"TargetAttrInterface"> {
meant to be used for passing additional options that are not in the
attribute.
}], "Attribute", "createObject",
- (ins "const SmallVector<char, 0>&":$object,
+ (ins "::mlir::Operation*":$module,
+ "const SmallVector<char, 0>&":$object,
"const gpu::TargetOptions&":$options)>
];
}
diff --git a/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp b/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
index 1e7596e8cc4af..ba9df3d00b324 100644
--- a/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/ModuleToBinary.cpp
@@ -114,7 +114,8 @@ LogicalResult moduleSerializer(GPUModuleOp op,
return failure();
}
- Attribute object = target.createObject(*serializedModule, targetOptions);
+ Attribute object =
+ target.createObject(op, *serializedModule, targetOptions);
if (!object) {
op.emitError("An error happened while creating the object.");
return failure();
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index acb903aa37caf..21e246a020da4 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -49,7 +49,7 @@ class NVVMTargetAttrImpl
serializeToObject(Attribute attribute, Operation *module,
const gpu::TargetOptions &options) const;
- Attribute createObject(Attribute attribute,
+ Attribute createObject(Attribute attribute, Operation *module,
const SmallVector<char, 0> &object,
const gpu::TargetOptions &options) const;
};
@@ -593,7 +593,7 @@ NVVMTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
}
Attribute
-NVVMTargetAttrImpl::createObject(Attribute attribute,
+NVVMTargetAttrImpl::createObject(Attribute attribute, Operation *module,
const SmallVector<char, 0> &object,
const gpu::TargetOptions &options) const {
auto target = cast<NVVMTargetAttr>(attribute);
diff --git a/mlir/lib/Target/LLVM/ROCDL/Target.cpp b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
index cc13e5b7436ea..cfa8a7c621d71 100644
--- a/mlir/lib/Target/LLVM/ROCDL/Target.cpp
+++ b/mlir/lib/Target/LLVM/ROCDL/Target.cpp
@@ -62,7 +62,7 @@ class ROCDLTargetAttrImpl
serializeToObject(Attribute attribute, Operation *module,
const gpu::TargetOptions &options) const;
- Attribute createObject(Attribute attribute,
+ Attribute createObject(Attribute attribute, Operation *module,
const SmallVector<char, 0> &object,
const gpu::TargetOptions &options) const;
};
@@ -473,7 +473,7 @@ std::optional<SmallVector<char, 0>> ROCDLTargetAttrImpl::serializeToObject(
}
Attribute
-ROCDLTargetAttrImpl::createObject(Attribute attribute,
+ROCDLTargetAttrImpl::createObject(Attribute attribute, Operation *module,
const SmallVector<char, 0> &object,
const gpu::TargetOptions &options) const {
gpu::CompilationTarget format = options.getCompilationTarget();
diff --git a/mlir/lib/Target/SPIRV/Target.cpp b/mlir/lib/Target/SPIRV/Target.cpp
index 4c416abe71cac..d48548bf9709c 100644
--- a/mlir/lib/Target/SPIRV/Target.cpp
+++ b/mlir/lib/Target/SPIRV/Target.cpp
@@ -34,7 +34,7 @@ class SPIRVTargetAttrImpl
serializeToObject(Attribute attribute, Operation *module,
const gpu::TargetOptions &options) const;
- Attribute createObject(Attribute attribute,
+ Attribute createObject(Attribute attribute, Operation *module,
const SmallVector<char, 0> &object,
const gpu::TargetOptions &options) const;
};
@@ -89,7 +89,7 @@ std::optional<SmallVector<char, 0>> SPIRVTargetAttrImpl::serializeToObject(
// Prepare Attribute for gpu.binary with serialized kernel object
Attribute
-SPIRVTargetAttrImpl::createObject(Attribute attribute,
+SPIRVTargetAttrImpl::createObject(Attribute attribute, Operation *module,
const SmallVector<char, 0> &object,
const gpu::TargetOptions &options) const {
gpu::CompilationTarget format = options.getCompilationTarget();
|
|
Can this be tested? |
As it stands, not really as this patch is just adding an argument that none of the upstream target attributes is using. |
fabianmcg
force-pushed
the
pr-gpu-add-module
branch
from
596875c to
8527ee1
Compare
|
@joker-eph to test it I decided to pass the discardable attributes in the gpu module to the |
|
Ping for review. |
joker-eph
reviewed
Jun 12, 2024
mlir/lib/Target/LLVM/NVVM/Target.cpp
Outdated
| if (format == gpu::CompilationTarget::Assembly) { | ||
| SmallVector<NamedAttribute, 1> attrs(objectProps.getValue()); | ||
| attrs.push_back( | ||
| builder.getNamedAttr("O", builder.getI32IntegerAttr(target.getO()))); |
There was a problem hiding this comment.
What if O was already in the dictionary?
That seems... surprising to me, and kind of arbitrary. Who's consuming this then? Why isn't it on the target attr in the first place? |
That wasn't their default behavior, I was making it so the approach was testable, because there are not that many ways to test this PR. However, with #95292 this PR will be testable with the ROCDL attribute. |
|
#95292 seems fine to me because as far as I understand it is selectively adding some relevant informations to the metadata instead of just being a pass-through. |
Sure, I'll drop the pass through. |
fabianmcg
force-pushed
the
pr-gpu-add-module
branch
from
8527ee1 to
5d7a281
Compare
|
Ping for review. |
This patch adds an argument to `gpu::TargetAttrInterface::createObject` to pass the GPU module. This is useful as `gpu::ObjectAttr` contains a property dict for metadata, and passing the module allows extracting things like the symbol table and adding it to the property dict.
fabianmcg
force-pushed
the
pr-gpu-add-module
branch
from
5d7a281 to
1de0959
Compare
|
Ping for review. |
ftynse
approved these changes
Aug 27, 2024
Co-authored-by: Oleksandr "Alex" Zinenko <ftynse@gmail.com>
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/116/builds/2735 Here is the relevant piece of the build log for the reference |
|
The failure is: It might be a |
I had issues with this bot before. It just randomly timed out in my cases. Might be some OOM issues? |
Same, it also has weird c++ quirks (g++-7). Fortunately the last build pass. |
When you see "Killed" it is an indication of OOM: the machine does not have enough RAM for the amount of parallel process running here. |
joker-eph
reviewed
Aug 27, 2024
| @@ -47,8 +47,9 @@ def GPUTargetAttrInterface : AttrInterface<"TargetAttrInterface"> { | |||
| meant to be used for passing additional options that are not in the | |||
| attribute. | |||
| }], "::mlir::Attribute", "createObject", | |||
| (ins "const ::llvm::SmallVector<char, 0>&":$object, | |||
| "const ::mlir::gpu::TargetOptions&":$options)> | |||
| (ins "::mlir::Operation *":$module, | |||
There was a problem hiding this comment.
We're missing documentation here!
(also should it be GPUModuleOp instead of Operation*)
There was a problem hiding this comment.
The GPUModuleOp vs Operation* goes back to the original introduction of these attributes. If you see, all other methods also use Operation*. Back then, there was a build dependency that required the interfaces to be generated before the dialect, hence the types were not available (I'm not sure this is needed anymore, I'll check). Then a second reason, was to make the header independent of GPUDialect.h
joker-eph
reviewed
Aug 27, 2024
| StringAttr::get(module->getContext(), | ||
| StringRef(object.data(), object.size())), | ||
| module->getAttrDictionary()); | ||
| } |
There was a problem hiding this comment.
Can you also add more documentation to the tests, out of the context of this PR it wouldn't be straightforward to understand the intent of all this I believe.
There was a problem hiding this comment.
Ok, wrt docs everywhere, would you prefer a fresh patch, or can I bundle the changes into this patch #95292 ? (they are related)
There was a problem hiding this comment.
mlir/include/mlir/Dialect/GPU/IR/CompilationAttrInterfaces.td isn't touched at all in #95292 ? I wouldn't add it to the changeset there.
There was a problem hiding this comment.
You're right, I'm touching CompilationAttrs.td. I'll merge that other one, create a patch with the docs and then revisit if we can change Operation* to GPUModuleOp.
fabianmcg
added a commit
to fabianmcg/llvm-project
that referenced
this pull request
Aug 27, 2024
Fix docs modified by llvm#94910 by adding information about the `module` argument in `gpu::TargetAttrInterface::createObject`.
fabianmcg
added a commit
that referenced
this pull request
Aug 27, 2024
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
This patch adds an argument to
gpu::TargetAttrInterface::createObjectto pass the GPU module. This is useful asgpu::ObjectAttrcontains a property dict for metadata, hence the module can be used for extracting things like the symbol table and adding it to the property dict.By default target attributes (NVVM, ROCDL, SPIRV) will pass all discardable attributes to the property dictionary.