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

[Antonio Frighetto via flang-commits]

================
@@ -1475,6 +1475,23 @@ let Predicates = [SM80] in {
             [(set i32:$res, (int_nvvm_mbarrier_pending_count i64:$state))]>;
 }
 
+let Predicates = [PTX93, SM90] in {
+  class MBARRIER_INIT_LAYOUT<NVPTXAddressSpace as> :
+            BasicNVPTXInst<(outs), (ins ADDR:$addr, B32:$count),
+            "mbarrier.init.layout::v1" # as.Suffix # ".b64",
+            [(IntrinsicInAS<int_nvvm_mbarrier_init, as>
+                  addr:$addr, i32:$count, (i32 1))]>;
+
+  def MBARRIER_INIT_LAYOUT : MBARRIER_INIT_LAYOUT<AddrSpaceGeneric>;
+  def MBARRIER_INIT_LAYOUT_SHARED : MBARRIER_INIT_LAYOUT<AddrSpaceShared>;
+
+  def MBARRIER_CHECK_LAYOUT_SHARED :
+      NVPTXInst<(outs B1:$res), (ins ADDR:$addr, B32:$layout),
+                "mbarrier.check_layout.layout::v${layout}.shared.b64 $res, [$addr];",
+                [(set i1:$res,
+                  (int_nvvm_mbarrier_check_layout addr:$addr, i32:$layout))]>;
----------------
antoniofrighetto wrote:

Just wondering, shouldn't the `int_nvvm_mbarrier_check_layout` declaration be constrained to `llvm_shared_ptr_ty` for the time being? IIUC from the ISA, it seems like both generic and .shared::cta should be supported for mbarrier.check_layout. If we intend to leave the address-space parameterization (as mbarrier.init does) as a follow-up, we might want to constrain this in the meantime (I guess we could reach ISel with no matching pattern).

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