-
Notifications
You must be signed in to change notification settings - Fork 11.9k
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
[mlir][GPU] Add builders to allow passing in integer upper_bound
s
#114252
Conversation
These are convendience builders to allow user code that knows (or optionally knows) the upper_bound on, say, a `gpu.thread_id` to specify that bound without needing to manually construct an IndexAttr.
@llvm/pr-subscribers-mlir Author: Krzysztof Drewniak (krzysz00) ChangesThese are convendience builders to allow user code that knows (or optionally knows) the upper_bound on, say, a Full diff: https://github.com/llvm/llvm-project/pull/114252.diff 1 Files Affected:
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index 6098eb34d04d52..f620431ca92ac9 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -74,7 +74,20 @@ class GPU_IndexOp<string mnemonic, list<Trait> traits = []> :
}]>,
OpBuilder<(ins "::mlir::Type":$resultType, "::mlir::gpu::Dimension":$dimension), [{
build($_builder, $_state, resultType, dimension, /*upperBound=*/nullptr);
+ }]>,
+ OpBuilder<(ins "::mlir::gpu::Dimension":$dimension, "std::optional<int64_t>":$upperBound), [{
+ ::mlir::IntegerAttr upperBoundAttr = nullptr;
+ if (upperBound.has_value())
+ upperBoundAttr = $_builder.getIndexAttr(*upperBound);
+ build($_builder, $_state, dimension, upperBoundAttr);
+ }]>,
+ OpBuilder<(ins "::mlir::Type":$resultType, "::mlir::gpu::Dimension":$dimension, "std::optional<int64_t>":$upperBound), [{
+ ::mlir::IntegerAttr upperBoundAttr = nullptr;
+ if (upperBound.has_value())
+ upperBoundAttr = $_builder.getIndexAttr(*upperBound);
+ build($_builder, $_state, resultType, dimension, upperBoundAttr);
}]>
+
];
}
|
@llvm/pr-subscribers-mlir-gpu Author: Krzysztof Drewniak (krzysz00) ChangesThese are convendience builders to allow user code that knows (or optionally knows) the upper_bound on, say, a Full diff: https://github.com/llvm/llvm-project/pull/114252.diff 1 Files Affected:
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index 6098eb34d04d52..f620431ca92ac9 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -74,7 +74,20 @@ class GPU_IndexOp<string mnemonic, list<Trait> traits = []> :
}]>,
OpBuilder<(ins "::mlir::Type":$resultType, "::mlir::gpu::Dimension":$dimension), [{
build($_builder, $_state, resultType, dimension, /*upperBound=*/nullptr);
+ }]>,
+ OpBuilder<(ins "::mlir::gpu::Dimension":$dimension, "std::optional<int64_t>":$upperBound), [{
+ ::mlir::IntegerAttr upperBoundAttr = nullptr;
+ if (upperBound.has_value())
+ upperBoundAttr = $_builder.getIndexAttr(*upperBound);
+ build($_builder, $_state, dimension, upperBoundAttr);
+ }]>,
+ OpBuilder<(ins "::mlir::Type":$resultType, "::mlir::gpu::Dimension":$dimension, "std::optional<int64_t>":$upperBound), [{
+ ::mlir::IntegerAttr upperBoundAttr = nullptr;
+ if (upperBound.has_value())
+ upperBoundAttr = $_builder.getIndexAttr(*upperBound);
+ build($_builder, $_state, resultType, dimension, upperBoundAttr);
}]>
+
];
}
|
@@ -74,7 +74,20 @@ class GPU_IndexOp<string mnemonic, list<Trait> traits = []> : | |||
}]>, | |||
OpBuilder<(ins "::mlir::Type":$resultType, "::mlir::gpu::Dimension":$dimension), [{ | |||
build($_builder, $_state, resultType, dimension, /*upperBound=*/nullptr); | |||
}]>, | |||
OpBuilder<(ins "::mlir::gpu::Dimension":$dimension, "std::optional<int64_t>":$upperBound), [{ |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Why is it optional?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
So that you can explicitly pass in std::nullopt
to leave the bound unset
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I also don't get what convenience does it bring. Why would you demand a nullopt when there can be two builders: with and without the attribute/whatever? Or if we want a single builder - a default value.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
upper_bound
is a fundamentally optional property.
This is designed to make people's lives easier when they're in possession of something like an ArrayRef<int64_t> blockSizes
.
I'll note we already have the attribute-less builder.
The advantage to having this be optional is that there's downstream code that either optionally knows the block sizes or optionally wants to forget about them.
I'm not sure I understand the motivation: is this just hiding the index attribute creation? In general I am not fond of builder which hide this kind of cost, even though I realize the tradeoff with usability here. Taking another angle: can we instead move to a native integer property and avoid the cost of the attribute creation? |
@joker-eph Yeah, you're right, refactoring to an |
Closing PR in favor of a properties refactor |
These are convendience builders to allow user code that knows (or optionally knows) the upper_bound on, say, a
gpu.thread_id
to specify that bound without needing to manually construct an IndexAttr.