================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
