[Mlir-commits] [llvm] [mlir] [LLVM][NVPTX][MLIR] Add mbarrier layout support (PR #217252)

[Guray Ozen]

================
@@ -704,6 +705,33 @@ def NVVM_MBarrierInvalOp : NVVM_VoidIntrinsicOp<"mbarrier.inval">,
   let assemblyFormat = "$addr attr-dict `:` type(operands)";
 }
 
+def NVVM_MBarrierCheckLayoutOp :
+    NVVM_SingleResultIntrinsicOp<"mbarrier.check_layout"> {
+  let summary = "MBarrier Check-Layout Operation";
+  let description = [{
+    The `nvvm.mbarrier.check_layout` operation tests whether the *mbarrier
+    object* at `addr` was initialized with the layout named by `layout`.
+
+    - `res`: An `i1` that is `true` when the *mbarrier object* has the queried
+      layout and `false` otherwise.
+
+    The operation takes the following operand and attribute:
+    - `addr`: A pointer to the memory location of the *mbarrier object*. The
+      `addr` must be a pointer to shared::cta memory.
+    - `layout`: The mbarrier layout version to test for. Only `0` and `1` are
+      valid values.
+
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-check-layout)
+  }];
+
+  let results = (outs I1:$res);
+  let arguments = (ins
+    LLVM_PointerShared:$addr,
+    ConfinedAttr<I32Attr, [IntMinValue<0>, IntMaxValue<1>]>:$layout);
----------------
grypp wrote:

Ditto - can this be default valued attribute here

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