[mlir][GPU] Add builders to allow passing in integer upper_bounds

[krzysz00]

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.

[llvmbot]

@llvm/pr-subscribers-mlir

Author: Krzysztof Drewniak (krzysz00)

Changes

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.

---

Full diff: https://github.com/llvm/llvm-project/pull/114252.diff

1 Files Affected:

• (modified) mlir/include/mlir/Dialect/GPU/IR/GPUOps.td (+13)

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);
     }]>
+
   ];
 }
 

[llvmbot]

@llvm/pr-subscribers-mlir-gpu

Author: Krzysztof Drewniak (krzysz00)

Changes

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.

---

Full diff: https://github.com/llvm/llvm-project/pull/114252.diff

1 Files Affected:

• (modified) mlir/include/mlir/Dialect/GPU/IR/GPUOps.td (+13)

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);
     }]>
+
   ];
 }
 

[kurapov-peter]

Why is it optional?

[krzysz00]

So that you can explicitly pass in std::nullopt to leave the bound unset

[kurapov-peter]

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.

[krzysz00]

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.

[joker-eph]

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?

[krzysz00]

@joker-eph Yeah, you're right, refactoring to an OptionalProperty<I64Property> is likely a better move here

[krzysz00]

Closing PR in favor of a properties refactor