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

[Krzysztof Drewniak]

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

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.

>From abc7b42fbcf3a112762fa55e97a9a5d2832f60cf Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Wed, 30 Oct 2024 15:43:35 +0000
Subject: [PATCH] [mlir][GPU] Add builders to allow passing in integer
 `upper_bound`s

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.
---
 mlir/include/mlir/Dialect/GPU/IR/GPUOps.td | 13 +++++++++++++
 1 file changed, 13 insertions(+)

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