[Mlir-commits] [mlir] [mlir][GPU] Add builders to allow passing in integer `upper_bound`s (PR #114252)

[llvmlistbot at llvm.org]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir

Author: Krzysztof Drewniak (krzysz00)

<details>
<summary>Changes</summary>

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

``````````

</details>


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