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

[Rajat Bajpai]

================
@@ -2157,9 +2157,23 @@ foreach is_shared = [true, false] in {
   }
 }
 
+def int_nvvm_mbarrier_init_layout :
+  DefaultAttrsIntrinsic<[], [llvm_shared_ptr_ty,  // mbar
+                              llvm_i32_ty,        // count
+                              llvm_i32_ty],       // layout
----------------
rajatbajpai wrote:

Or we could do the following:

1. Keep mbarrier.init (without the layout in the name), add an extra argument for the layout, and then handle it through AU.
2. With this approach, it is important that we continue generating the default PTX mbarrier instruction. We don't want to AutoUpgrade older mbarrier uses to the newer form with layout modifiers.

mbarrier.init (layout = 0) => mbarrier.init{.shared{::cta}}.b64
mbarrier.init (layout = 1) => mbarrier.init.layout::v1{.shared{::cta}}.b64

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