[Mlir-commits] [mlir] [MLIR][NVGPU] Add `mbarrier.get` Op (PR #133221)

[Guray Ozen]

================
@@ -322,6 +322,25 @@ def NVGPU_MBarrierCreateOp : NVGPU_Op<"mbarrier.create", []> {
   }];
 }
 
+def NVGPU_MBarrierGetOp : NVGPU_Op<"mbarrier.get", []> {
+  let summary = "Return a pointer to an `nvgpu.mbarrier`.";
+  let description = [{
+    The `nvgpu.mbarrier.get` operation retrieves a pointer to a specific 
+    `mbarrier` object from a group of barriers created by the `nvgpu.mbarrier.create` operation.
+
+    Example:
+    ```mlir
+      %mbars = nvgpu.mbarrier.create -> !nvgpu.mbarrier.group<memorySpace = #gpu.address_space<workgroup>, num_barriers = 10>
+      %mbar_pointer = nvgpu.mbarrier.get %mbars[%c2] : !nvgpu.mbarrier.group<memorySpace = #gpu.address_space<workgroup>>
----------------
grypp wrote:

We can only check locals to the op in the verifier. Here `idx` is an SSA value, and it requires some data flow analysis. According the guide, we should not do that:
https://mlir.llvm.org/getting_started/DeveloperGuide/#ir-verifier

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