[Mlir-commits] [llvm] [mlir] [NVPTX] Add "exclusive" intrinsic variants for tcgen05.alloc/dealloc (PR #216016)

[Dharuni R Acharya]

================
@@ -5735,30 +5735,85 @@ let Predicates = [SM90] in {
 def EXIT : NullaryInst<"exit", int_nvvm_exit>;
 
 // Tcgen05 intrinsics
-let isConvergent = true in {
-let Predicates = [hasTcgen05InstSupport] in {
-multiclass TCGEN05_ALLOC_INTR<string AS, string num, Intrinsic Intr> {
-  def "" : BasicNVPTXInst<(outs),
-             (ins ADDR:$dst, B32:$ncols),
-             "tcgen05.alloc.cta_group::" # num # ".sync.aligned" # AS # ".b32",
-             [(Intr addr:$dst, B32:$ncols)]>;
+multiclass TCGEN05_ALLOC_PATFRAGS<Intrinsic Intr> {
+  defvar frag_pat = (Intr node:$dst, node:$ncols);
+  def _GENERIC
+      : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.generic>;
+  def _SHARED
+      : PatFrag<!setdagop(frag_pat, ops), frag_pat, AS_match.shared>;
+}
----------------
DharuniRAcharya wrote:

Makes sense! Updated in the latest revision! 
Thanks!

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