[llvm] [mlir] [LLVM][NVPTX][MLIR] Add mbarrier layout and Rubin multicast extensions (PR #216753)

[Rajat Bajpai via llvm-commits]

================
@@ -635,7 +635,10 @@ def NVVM_PMEventOp : NVVM_VoidIntrinsicOp<"pmevent">,
 /// mbarrier.init instruction with generic pointer type
 def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
   Arguments<(ins AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr,
-                 I32:$count, PtxPredicate:$predicate)> {
+                 I32:$count,
----------------
rajatbajpai wrote:

@grypp @schwarzschild-radius I think we can have two logical buckets:
1. New mbarrier layout support in mbarrier.
2. 32-bit multicast mask support in mbarriers.

I think it would be easier to review if init and check layout goes together.

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