[llvm] [mlir] [LLVM][NVPTX] Add async bulk copy global to shared extensions (PR #222323)

[Durgadoss R via llvm-commits]

================
@@ -1264,6 +1264,55 @@ static Intrinsic::ID shouldUpgradeNVPTXTMAG2SCTAIntrinsics(Function *F,
 
   return ID;
 }
+// The legacy tail of llvm.nvvm.cp.async.bulk.global.to.shared.cluster is:
+//
+//   ..., i16 %mc, i64 %ch, i1 %flag_mc, i1 %flag_ch
+//
+// The current intrinsic is overloaded on the multicast-mask type and takes a
+// trailing i32 %validate_pattern; the legacy tail is recognized by an i1 at
+// parameter N-1.
+static Intrinsic::ID
+shouldUpgradeNVPTXBulkG2SClusterIntrinsic(Function *F, StringRef Name,
+                                          SmallVectorImpl<Type *> &OvlTys) {
+  if (!Name.consume_front("cp.async.bulk.global.to.shared.cluster"))
+    return Intrinsic::not_intrinsic;
+
+  // Parameter N-1 is i1 for the legacy tail; the current tail ends with
+  // i32 %validate_pattern, for which N-1 is i32.
+  size_t NumParams = F->getFunctionType()->getNumParams();
+  if (!F->getFunctionType()->getParamType(NumParams - 1)->isIntegerTy(1))
+    return Intrinsic::not_intrinsic;
+
+  // The multicast mask is parameter 4; legacy IR only uses i16.
+  Type *MaskTy = F->getFunctionType()->getParamType(NumParams - 4);
+  if (!MaskTy->isIntegerTy(16))
+    return Intrinsic::not_intrinsic;
+  OvlTys.push_back(MaskTy);
+
+  return Intrinsic::nvvm_cp_async_bulk_global_to_shared_cluster;
+}
+
+// The legacy tail of llvm.nvvm.cp.async.bulk.global.to.shared.cta is:
+//
+//   ..., i64 %ch, i1 %flag_ch
+//
+// The current intrinsic adds %ignore_bytes_left/%ignore_bytes_right before
+// %ch and trailing %flag_oob/%validate_pattern; the legacy tail is
+// recognized by an i1 at parameter N-1, whereas the current tail ends
+// with an i32.
+static Intrinsic::ID shouldUpgradeNVPTXBulkG2SCTAIntrinsic(Function *F,
+                                                           StringRef Name) {
----------------
durga4github wrote:

ok, I missed the fact that ignore_bytes_left/right are not immediates. Let us stick to our convention of values followed by flags. So, we can keep this as is.
+
On the other question of doing these as an NFC later, I am fine with that.

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