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

[Rajat Bajpai via llvm-commits]

================
@@ -524,20 +524,35 @@ let Predicates = [PTX80, SM90] in {
 // TMA Async Bulk Copy Functions
 //------------------------------
 
-class CpAsyncBulkStr<bit mc, bit ch, bit mask = 0> {
+class CpAsyncBulkStr<bits<2> mc, bit ch, bit mask = 0, bit is_relaxed = 0,
+                     bit oob = 0> {
+  // Multicast kind: 0 = none, 1 = 16-bit mask, 2 = 32-bit mask.
+  string mc_str = !cond(!eq(mc, 1): ".multicast::cluster",
+                        !eq(mc, 2): ".multicast::cluster::32b",
+                        true: "");
+
   // Shared to Global memory
   string S2G = "cp.async.bulk.global.shared::cta.bulk_group"
                # !if(ch, ".L2::cache_hint", "")
                # !if(mask, ".cp_mask", "");
 
-  // Global to Shared cluster memory
-  string G2S = "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes"
-               # !if(mc, ".multicast::cluster", "")
-               # !if(ch, ".L2::cache_hint", "");
+  // Global to Shared cluster memory. The relaxed form carries the "${scope}"
+  // qualifier and an explicit ".b128" suffix; the weak form has neither.
+  defvar prefix = !if(is_relaxed,
+      "cp.async.bulk.relaxed${scope}.shared::cluster.global.mbarrier::complete_tx::bytes$validate",
+      "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes$validate");
+  string G2S = prefix # mc_str
+               # !if(ch, ".L2::cache_hint", "")
+               # !if(is_relaxed, ".b128", "");
----------------
rajatbajpai wrote:

done.

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